--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/src/Atomic.h	Thu Jan 16 15:12:06 2020 +0900
@@ -0,0 +1,9 @@
+typedef struct Atomic<Impl>{
+    union Data* atomic;
+    union Data** ptr;
+    union Data* oldData;
+    union Data* newData;
+    __code checkAndSet(Impl* atomic, union Data** ptr, union Data* oldData, union Data* newData, __code next(...), __code fail(...));
+    __code next(...);
+    __code fail(...);
+} Atomic;

--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/src/AtomicReference.cbc	Thu Jan 16 15:12:06 2020 +0900
@@ -0,0 +1,18 @@
+#include "../context.h"
+#interface "Atomic.h"
+#include <stdio.h>
+
+Atomic* createAtomicReference(struct Context* context) {
+    struct Atomic* atomic = new Atomic();
+    struct AtomicReference* atomicReference = new AtomicReference();
+    atomic->atomic = (union Data*)atomicReference;
+    atomic->checkAndSet = C_checkAndSetAtomicReference;
+    return atomic;
+}
+
+__code checkAndSetAtomicReference(struct AtomicReference* atomic, union Data** ptr, union Data* oldData, union Data* newData, __code next(...), __code fail(...)) {
+    if (__sync_bool_compare_and_swap(ptr, oldData, newData)) {
+        goto next(...);
+    }
+    goto fail(...);
+}

--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/src/Buffer.h	Thu Jan 16 15:12:06 2020 +0900
@@ -0,0 +1,7 @@
+typedef struct Buffer<Impl>{
+        union Data* buffer;
+        union Data* data;
+        __code put(Impl* buffer, union Data* data, __code next(...));
+        __code take(Impl* buffer, __code next(union Data*, ...));
+        __code next(...);
+} Buffer;

--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/src/CMakeLists.txt	Thu Jan 16 15:12:06 2020 +0900
@@ -0,0 +1,142 @@
+cmake_minimum_required(VERSION 3.8)
+
+set(USE_CUDA,0)
+# -DUSE_CUDA
+#  add_definitions("-Wall -g -O")
+
+set(CMAKE_C_COMPILER $ENV{CBC_COMPILER})
+add_definitions("-Wall -g")
+
+# -DCMAKE_BUILD_TYPE=Debug
+set(CMAKE_C_FLAGS_DEBUG "-O0")
+
+if (${USE_CUDA})
+    include_directories("/usr/local/cuda/include")
+    set(NVCCFLAG "-std=c++11" "-g" "-O0" )
+    if (UNIX AND NOT APPLE) # LINUX
+        set(CUDA_LINK_FLAGS "-L/usr/local/cuda/lib64 -lcuda -lcudart")
+    elseif (APPLE)
+        set(CUDA_LINK_FLAGS "-framework CUDA -lc++ -Wl,-search_paths_first -Wl,-headerpad_max_install_names /usr/local/cuda/lib/libcudart_static.a -Wl,-rpath,/usr/local/cuda/lib")
+    endif()
+    find_package(CUDA REQUIRED)
+    SET( CMAKE_EXE_LINKER_FLAGS  "${CMAKE_EXE_LINKER_FLAGS} ${CUDA_LINK_FLAGS}" )
+endif()
+
+macro( GearsCommand )
+    set( _OPTIONS_ARGS )
+    set( _ONE_VALUE_ARGS TARGET )
+    set( _MULTI_VALUE_ARGS SOURCES )
+    cmake_parse_arguments( _Gears "${_OPTIONS_ARGS}" "${_ONE_VALUE_ARGS}" "${_MULTI_VALUE_ARGS}" ${ARGN} )
+
+    set (_Gears_CSOURCES)
+    foreach(i ${_Gears_SOURCES})
+        if (${i} MATCHES "\\.cbc")
+            string(REGEX REPLACE "(.*).cbc" "c/\\1.c" j ${i})
+            add_custom_command (
+                OUTPUT    ${j}
+                DEPENDS   ${i}
+                COMMAND  "perl" "generate_stub.pl" "-o" ${j} ${i}
+            )
+        elseif (${i} MATCHES "\\.cu")
+            string(REGEX REPLACE "(.*).cu" "c/\\1.ptx" j ${i})
+            add_custom_command (
+                OUTPUT    ${j}
+                DEPENDS   ${i}
+                COMMAND  nvcc ${NVCCFLAG} -c -ptx -o ${j} ${i}
+            )
+        else()
+            set(j ${i})
+        endif()
+        list(APPEND _Gears_CSOURCES ${j})
+    endforeach(i)
+
+    add_custom_command (
+          OUTPUT    c/${_Gears_TARGET}-context.c
+          DEPENDS   ${_Gears_CSOURCES}
+          COMMAND  "perl" "generate_context.pl" "-o" ${_Gears_TARGET} "-w" ${_Gears_CSOURCES}
+    )
+    add_executable(${_Gears_TARGET} ${_Gears_CSOURCES} c/${_Gears_TARGET}-context.c)
+    target_link_libraries(${_Gears_TARGET} m pthread)
+endmacro()
+
+
+GearsCommand(
+  TARGET
+      twice
+  SOURCES
+      examples/twice/main.cbc examples/twice/createArray.cbc examples/twice/twice.cbc examples/twice/printArray.cbc CPUWorker.cbc TaskManagerImpl.cbc SingleLinkedQueue.cbc SynchronizedQueue.cbc TimerImpl.cbc MultiDimIterator.cbc AtomicReference.cbc
+)
+
+GearsCommand(
+  TARGET
+      calc
+  SOURCES
+      examples/calc/calc.cbc examples/calc/add.cbc examples/calc/mult.cbc examples/calc/initIntegerDataGears.cbc CPUWorker.cbc TaskManagerImpl.cbc SingleLinkedQueue.cbc SynchronizedQueue.cbc AtomicReference.cbc
+)
+
+GearsCommand(
+  TARGET
+      bitonicSort
+  SOURCES
+      examples/bitonicSort/bitonicSort.cbc examples/bitonicSort/bitonicSwap.cbc examples/bitonicSort/makeArray.cbc examples/bitonicSort/printArray.cbc CPUWorker.cbc TaskManagerImpl.cbc SingleLinkedQueue.cbc SynchronizedQueue.cbc MultiDimIterator.cbc TimerImpl.cbc AtomicReference.cbc
+)
+
+if (${USE_CUDA})
+    GearsCommand(
+      TARGET
+          CUDAtwice
+      SOURCES 
+          examples/twice/main.cbc examples/twice/twice.cbc examples/twice/CUDAtwice.cu examples/twice/createArray.cbc examples/twice/printArray.cbc CPUWorker.cbc TimerImpl.cbc examples/twice/twice.cbc TaskManagerImpl.cbc SingleLinkedQueue.cbc SynchronizedQueue.cbc CUDAWorker.cbc cuda.c MultiDimIterator.cbc CUDAExecutor.cbc AtomicReference.cbc 
+    )
+    set_target_properties(CUDAtwice PROPERTIES COMPILE_FLAGS "-Wall -g -DUSE_CUDAWorker=1")
+
+    GearsCommand(
+      TARGET
+          CUDAbitonicSort
+      SOURCES 
+          examples/bitonicSort/bitonicSort.cbc examples/bitonicSort/bitonicSwap.cbc examples/bitonicSort/CUDAbitonicSwap.cu examples/bitonicSort/makeArray.cbc examples/bitonicSort/printArray.cbc CPUWorker.cbc CUDAWorker.cbc TaskManagerImpl.cbc SingleLinkedQueue.cbc SynchronizedQueue.cbc cuda.c MultiDimIterator.cbc TimerImpl.cbc CUDAExecutor.cbc AtomicReference.cbc
+    )
+    set_target_properties(CUDAbitonicSort PROPERTIES COMPILE_FLAGS "-Wall -g -DUSE_CUDAWorker=1")
+endif()
+
+GearsCommand(
+  TARGET
+      queue_test
+  SOURCES 
+      test/queue_test.cbc SingleLinkedQueue.cbc
+)
+
+GearsCommand(
+  TARGET
+      stack_test
+  SOURCES 
+      test/stack_test.cbc SingleLinkedStack.cbc SingleLinkedQueue.cbc
+)
+
+GearsCommand(
+  TARGET
+      multiDimIterator_test
+  SOURCES
+      test/multiDimIterator_test.cbc test/printIterator.cbc CPUWorker.cbc TaskManagerImpl.cbc SingleLinkedQueue.cbc SynchronizedQueue.cbc MultiDimIterator.cbc AtomicReference.cbc
+)
+
+#GearsCommand(
+#  TARGET
+#      sort
+#  SOURCES
+#      examples/bitonicSort/sort.cbc
+#)
+
+GearsCommand(
+  TARGET
+      rbtree
+  SOURCES
+      SingleLinkedQueue.cbc test/rbTree_test.cbc RedBlackTree.cbc SingleLinkedStack.cbc compare.c
+)
+
+GearsCommand(
+  TARGET
+      boundedBuffer
+  SOURCES
+  examples/boundedBuffer/main.cbc examples/boundedBuffer/initBuffer.cbc examples/boundedBuffer/SemaphoreImpl.cbc examples/boundedBuffer/BoundedBuffer.cbc examples/boundedBuffer/consumer.cbc examples/boundedBuffer/producer.cbc SpinLock.cbc CPUWorker.cbc TaskManagerImpl.cbc SingleLinkedQueue.cbc SynchronizedQueue.cbc MultiDimIterator.cbc AtomicReference.cbc
+)

--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/src/CPUWorker.cbc	Thu Jan 16 15:12:06 2020 +0900
@@ -0,0 +1,124 @@
+#include "../context.h"
+#interface "TaskManager.h"
+#interface "Worker.h"
+#interface "Iterator.h"
+#interface "Queue.h"
+
+static void startWorker(Worker* worker);
+
+Worker* createCPUWorker(struct Context* context, int id, Queue* queue) {
+    struct Worker* worker = new Worker();
+    struct CPUWorker* cpuWorker = new CPUWorker();
+    worker->worker = (union Data*)cpuWorker;
+    worker->tasks = queue;
+    cpuWorker->id = id;
+    cpuWorker->loopCounter = 0;
+    worker->taskReceive = C_taskReceiveCPUWorker;
+    worker->shutdown = C_shutdownCPUWorker;
+    pthread_create(&worker->thread, NULL, (void*)&startWorker, worker);
+    return worker;
+}
+
+static void startWorker(struct Worker* worker) {
+    struct CPUWorker* cpuWorker = &worker->worker->CPUWorker;
+    cpuWorker->context = NEW(struct Context);
+    initContext(cpuWorker->context);
+    Gearef(cpuWorker->context, Worker)->worker = (union Data*)worker;
+    Gearef(cpuWorker->context, Worker)->tasks = worker->tasks;
+    goto meta(cpuWorker->context, worker->taskReceive);
+}
+
+__code taskReceiveCPUWorker(struct CPUWorker* worker, struct Queue* tasks) {
+    goto tasks->take(getTaskCPUWorker);
+}
+
+__code getTaskCPUWorker(struct CPUWorker* cpuWorker, struct Context* task, struct Worker* worker) {
+    if (!task) {
+        goto worker->shutdown(); // end thread
+    }
+    task->worker = worker;
+    enum Code taskCg = task->next;
+    task->next = C_odgCommitCPUWorker; // commit outputDG after task exec
+    goto meta(task, taskCg); // switch task context
+}
+
+__code getTaskCPUWorker_stub(struct Context* context) {
+    CPUWorker* cpuWorker = (CPUWorker*)GearImpl(context, Worker, worker);
+    Worker* worker = &Gearef(context,Worker)->worker->Worker;
+    struct Context* task = &Gearef(context, Queue)->data->Context;
+    goto getTaskCPUWorker(context, cpuWorker, task, worker);
+}
+
+__code odgCommitCPUWorker(struct CPUWorker* worker, struct Context* task) {
+    if (task->iterate) {
+        struct Iterator* iterator = task->iterator;
+        goto iterator->barrier(task, odgCommitCPUWorker1, odgCommitCPUWorker6);
+    } else {
+        goto odgCommitCPUWorker1();
+    }
+}
+
+__code odgCommitCPUWorker_stub(struct Context* context) {
+    // switch worker context
+    struct Context* workerContext = context->worker->worker->CPUWorker.context;
+    Gearef(workerContext, Worker)->worker = (union Data*)context->worker;
+    Gearef(workerContext, Worker)->task = context;
+    CPUWorker* cpuWorker = (CPUWorker*)GearImpl(workerContext, Worker, worker);
+    goto odgCommitCPUWorker(workerContext,
+                            cpuWorker,
+                            context);
+}
+
+__code odgCommitCPUWorker1(struct CPUWorker* worker, struct Context* task) {
+    int i = worker->loopCounter;
+    if (task->odg+i < task->maxOdg) {
+        goto odgCommitCPUWorker2();
+    }
+    worker->loopCounter = 0;
+    struct TaskManager* taskManager = task->taskManager;
+    goto taskManager->decrementTaskCount(odgCommitCPUWorker6);
+}
+
+__code odgCommitCPUWorker2(struct CPUWorker* worker, struct Context* task) {
+    int i = worker->loopCounter;
+    struct Queue* queue = GET_WAIT_LIST(task->data[task->odg+i]);
+    goto queue->isEmpty(odgCommitCPUWorker3, odgCommitCPUWorker5);
+}
+
+__code odgCommitCPUWorker3(struct CPUWorker* worker, struct Context* task) {
+    int i = worker->loopCounter;
+    struct Queue* queue = GET_WAIT_LIST(task->data[task->odg+i]);
+    goto queue->take(odgCommitCPUWorker4);
+}
+
+__code odgCommitCPUWorker4(struct CPUWorker* worker, struct Context* task, struct Context* waitTask) {
+    if (__sync_fetch_and_sub(&waitTask->idgCount, 1) == 1) { // atomic decrement idg counter(__sync_fetch_and_sub function return initial value of waitTask->idgCount point)
+        struct TaskManager* taskManager = waitTask->taskManager;
+        goto taskManager->spawn(waitTask, odgCommitCPUWorker2);
+    }
+    goto odgCommitCPUWorker2();
+}
+
+__code odgCommitCPUWorker4_stub(struct Context* context) {
+    CPUWorker* cpuWorker     = (CPUWorker*)GearImpl(context, Worker, worker);
+    struct Context* task     = Gearef(context, Worker)->task;
+    struct Context* waitTask = &Gearef(context, Queue)->data->Context;
+    goto odgCommitCPUWorker4(context,
+                             cpuWorker,
+                             task,
+                             waitTask);
+}
+
+__code odgCommitCPUWorker5(struct CPUWorker* worker, struct Context* task) {
+    worker->loopCounter++;
+    goto odgCommitCPUWorker1();
+}
+
+__code odgCommitCPUWorker6(struct CPUWorker* worker, struct Context* task) {
+    struct Worker* taskWorker = task->worker;
+    goto taskWorker->taskReceive(taskWorker->tasks);
+}
+
+__code shutdownCPUWorker(struct CPUWorker* worker) {
+    goto exit_code();
+}

--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/src/CUDAExecutor.cbc	Thu Jan 16 15:12:06 2020 +0900
@@ -0,0 +1,108 @@
+#include "../context.h"
+#interface "Executor.h"
+#interface "Timer.h"
+#include <stdio.h>
+#include <math.h>
+
+Executor* createCUDAExecutor(struct Context* context, CUdevice device) {
+    struct Executor* executor = new Executor();
+    struct CUDAExecutor* cudaExecutor = new CUDAExecutor();
+    cudaExecutor->timer = createTimerImpl(context);
+    checkCudaErrors(cuDeviceGetAttribute(&cudaExecutor->maxThreadPerBlock, CU_DEVICE_ATTRIBUTE_MAX_THREADS_PER_BLOCK, device));
+    executor->executor = (union Data*)cudaExecutor;
+    executor->read  = C_readCUDAExecutor;
+    executor->exec  = C_execCUDAExecutor;
+    executor->write = C_writeCUDAExecutor;
+    return executor;
+}
+
+__code readCUDAExecutor(struct CUDAExecutor* executor, struct Context* task, __code next(...)) {
+    struct CUDABuffer* buffer = executor->buffer;
+    int paramLen = buffer->inputLen + buffer->outputLen;
+    executor->kernelParams = (CUdeviceptr**)ALLOCATE_PTR_ARRAY(context, CUdeviceptr, paramLen);
+    for (int i = 0; i < paramLen; i++) {
+        CUdeviceptr* deviceptr = new CUdeviceptr();
+        // memory allocate
+        union Data* data = i < buffer->inputLen? buffer->inputData[i] : buffer->outputData[i-buffer->inputLen];
+        checkCudaErrors(cuMemAlloc(deviceptr, GET_SIZE(data)));
+        checkCudaErrors(cuMemcpyHtoD(*deviceptr, data, GET_SIZE(data)));
+        // Synchronous data transfer(host to device)
+        executor->kernelParams[i] = deviceptr;
+    }
+    // TODO: Implements pipeline
+    // goto next(...);
+    struct Timer* timer = executor->timer;
+    goto timer->start(execCUDAExecutor);
+}
+
+int computeblockDim(int count, int maxThreadPerBlock) {
+    return count < maxThreadPerBlock ? count : maxThreadPerBlock;
+}
+
+void calcBlockMaxThread(struct MultiDimIterator* iterator, struct CUDAExecutor* executor) {
+    executor->maxThreadPerBlockX = 1;
+    executor->maxThreadPerBlockY = 1;
+    executor->maxThreadPerBlockZ = 1;
+    // maxThreadPerBlockX * maxThreadPerBlockY * maxThreadPerBlockZ <= maxThreadPerBlock
+    if (iterator->x > 1 && iterator->y == 1 && iterator->z == 1) {
+        executor->maxThreadPerBlockX = executor->maxThreadPerBlock;
+        executor->maxThreadPerBlockY = 1;
+        executor->maxThreadPerBlockZ = 1;
+    } else if (iterator->x > 1 && iterator->y > 1 && iterator->z == 1) {
+        int ln_2 = log2(executor->maxThreadPerBlock);
+        int maxThread = 1 << (ln_2/2);
+        executor->maxThreadPerBlockX = maxThread;
+        executor->maxThreadPerBlockY = maxThread;
+        executor->maxThreadPerBlockZ = 1;
+    } else {
+        int ln_2 = log2(executor->maxThreadPerBlock);
+        int maxThread = 1 << (ln_2/3);
+        executor->maxThreadPerBlockX = maxThread * (1 << (ln_2%3));
+        executor->maxThreadPerBlockY = maxThread;
+        executor->maxThreadPerBlockZ = maxThread;
+    }
+}
+
+__code execCUDAExecutor(struct CUDAExecutor* executor, struct Context* task, __code next(...)) {
+    task->num_exec = 1;
+    if (task->iterate) {
+        struct MultiDimIterator* iterator = &task->iterator->iterator->MultiDimIterator;
+        calcBlockMaxThread(iterator, executor);
+        int blockDimX = computeblockDim(iterator->x, executor->maxThreadPerBlockX);
+        int blockDimY = computeblockDim(iterator->y, executor->maxThreadPerBlockY);
+        int blockDimZ = computeblockDim(iterator->z, executor->maxThreadPerBlockZ);
+        // launch kernel
+        checkCudaErrors(cuLaunchKernel(task->function,
+                    iterator->x/blockDimX, iterator->y/blockDimY, iterator->z/blockDimZ,
+                    blockDimX, blockDimY, blockDimZ,
+                    0, NULL, (void**)executor->kernelParams, NULL));
+    } else {
+        checkCudaErrors(cuLaunchKernel(task->function,
+                    1, 1, 1,
+                    1, 1, 1,
+                    0, NULL, (void**)executor->kernelParams, NULL));
+    }
+    // TODO: Implements pipeline
+    // goto next(...);
+    goto writeCUDAExecutor();
+}
+
+__code writeCUDAExecutor(struct CUDAExecutor* executor, struct Context* task, __code next(...)) {
+    // Asynchronous launch kernel
+    checkCudaErrors(cuCtxSynchronize());
+    struct Timer* timer = executor->timer;
+    goto timer->end(writeCUDAExecutor1);
+}
+
+__code writeCUDAExecutor1(struct CUDAExecutor* executor, struct Context* task, __code next(...)) {
+    //結果を取ってくるコマンドを入力する
+    struct CUDABuffer* buffer = executor->buffer;
+    int paramLen = buffer->inputLen + buffer->outputLen;
+    for (int i = 0; i < paramLen; i++) {
+        CUdeviceptr deviceptr =  *(executor->kernelParams[i]);
+        union Data* data = i < buffer->inputLen? buffer->inputData[i] : buffer->outputData[i-buffer->inputLen];
+        checkCudaErrors(cuMemcpyDtoH(data, deviceptr, GET_SIZE(data)));
+        cuMemFree(deviceptr);
+    }
+    goto next(...);
+}

--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/src/CUDAWorker.cbc	Thu Jan 16 15:12:06 2020 +0900
@@ -0,0 +1,131 @@
+#include "../context.h"
+#interface "TaskManager.h"
+#interface "Worker.h"
+#interface "Iterator.h"
+#interface "Queue.h"
+
+extern void cudaInit(struct CUDAWorker *cudaWorker,int phase, int deviceNum);
+extern void cudaShutdown(CUDAWorker *cudaWorker);
+
+static void startCUDAWorker(Worker* worker);
+
+Worker* createCUDAWorker(struct Context* context, int id, Queue* queue, int deviceNum) {
+    struct Worker* worker = new Worker();
+    struct CUDAWorker* cudaWorker = new CUDAWorker();
+    worker->worker = (union Data*)cudaWorker;
+    worker->tasks = queue;
+    cudaWorker->id = id;
+    cudaWorker->loopCounter = 0;
+    cudaWorker->deviceNum = deviceNum;
+    worker->taskReceive = C_taskReceiveCUDAWorker;
+    worker->shutdown = C_shutdownCUDAWorker;
+    pthread_create(&worker->thread, NULL, (void*)&startCUDAWorker, worker);
+    return worker;
+}
+
+static void startCUDAWorker(Worker* worker) {
+    struct CUDAWorker* cudaWorker = &worker->worker->CUDAWorker;
+    cudaInit(cudaWorker, 0, cudaWorker->deviceNum);
+    cudaWorker->context  = NEW(struct Context);
+    initContext(cudaWorker->context);
+    cudaWorker->executor = createCUDAExecutor(cudaWorker->context, cudaWorker->device);
+    Gearef(cudaWorker->context, Worker)->worker = (union Data*)worker;
+    Gearef(cudaWorker->context, Worker)->tasks = worker->tasks;
+    goto meta(cudaWorker->context, worker->taskReceive);
+}
+
+__code taskReceiveCUDAWorker(struct Worker* worker, struct Queue* tasks) {
+    goto tasks->take(getTaskCUDAWorker);
+}
+
+__code getTaskCUDAWorker(struct CUDAWorker* cudaWorker, struct Context* task, struct Worker* worker) {
+    if (!task) {
+        goto worker->shutdown(); // end thread
+    }
+    task->worker = worker;
+    enum Code taskCg = task->next;
+    task->next = C_odgCommitCUDAWorker; // commit outputDG after task exec
+    goto meta(task, taskCg); // switch task context
+}
+
+__code getTaskCUDAWorker_stub(struct Context* context) {
+    CUDAWorker* cudaWorker = (CUDAWorker*)GearImpl(context, Worker, worker);
+    Worker* worker = &Gearef(context,Worker)->worker->Worker;
+    struct Context* task = &Gearef(context, Queue)->data->Context;
+    goto getTaskCUDAWorker(context, cudaWorker, task, worker);
+}
+
+__code odgCommitCUDAWorker(struct CUDAWorker* worker, struct Context* task) {
+    if (task->iterate) {
+        struct Iterator* iterator = task->iterator;
+        goto iterator->barrier(task, odgCommitCUDAWorker1, odgCommitCUDAWorker6);
+    } else {
+        goto odgCommitCUDAWorker1();
+    }
+}
+
+__code odgCommitCUDAWorker_stub(struct Context* context) {
+    // switch worker context
+    struct Context* workerContext = context->worker->worker->CUDAWorker.context;
+    Gearef(workerContext, Worker)->worker = (union Data*)context->worker;
+    Gearef(workerContext, Worker)->task = context;
+    CUDAWorker* cudaWorker = (CUDAWorker*)GearImpl(workerContext, Worker, worker);
+    goto odgCommitCUDAWorker(workerContext,
+                            cudaWorker,
+                            context);
+}
+
+__code odgCommitCUDAWorker1(struct CUDAWorker* worker, struct Context* task) {
+    int i = worker->loopCounter;
+    if (task->odg+i < task->maxOdg) {
+        goto odgCommitCUDAWorker2();
+    }
+    worker->loopCounter = 0;
+    struct TaskManager* taskManager = task->taskManager;
+    goto taskManager->decrementTaskCount(odgCommitCUDAWorker6);
+}
+
+__code odgCommitCUDAWorker2(struct CUDAWorker* worker, struct Context* task) {
+    int i = worker->loopCounter;
+    struct Queue* queue = GET_WAIT_LIST(task->data[task->odg+i]);
+    goto queue->isEmpty(odgCommitCUDAWorker3, odgCommitCUDAWorker5);
+}
+
+__code odgCommitCUDAWorker3(struct CUDAWorker* worker, struct Context* task) {
+    int i = worker->loopCounter;
+    struct Queue* queue = GET_WAIT_LIST(task->data[task->odg+i]);
+    goto queue->take(odgCommitCUDAWorker4);
+}
+
+__code odgCommitCUDAWorker4(struct CUDAWorker* worker, struct Context* task, struct Context* waitTask) {
+    if (__sync_fetch_and_sub(&waitTask->idgCount, 1) == 1) { // atomic decrement idg counter(__sync_fetch_and_sub function return initial value of waitTask->idgCount point)
+        struct TaskManager* taskManager = waitTask->taskManager;
+        goto taskManager->spawn(waitTask, odgCommitCUDAWorker2);
+    }
+    goto odgCommitCUDAWorker2();
+}
+
+__code odgCommitCUDAWorker4_stub(struct Context* context) {
+    CUDAWorker* cudaWorker     = (CUDAWorker*)GearImpl(context, Worker, worker);
+    struct Context* task     = Gearef(context, Worker)->task;
+    struct Context* waitTask = &Gearef(context, Queue)->data->Context;
+    goto odgCommitCUDAWorker4(context,
+                             cudaWorker,
+                             task,
+                             waitTask);
+}
+
+__code odgCommitCUDAWorker5(struct CUDAWorker* worker, struct Context* task) {
+    worker->loopCounter++;
+    goto odgCommitCUDAWorker1();
+}
+
+__code odgCommitCUDAWorker6(struct CUDAWorker* worker, struct Context* task) {
+    struct Worker* taskWorker = task->worker;
+    goto taskWorker->taskReceive(taskWorker->tasks);
+}
+
+__code shutdownCUDAWorker(struct CUDAWorker* worker) {
+    cudaShutdown(worker);
+    goto meta(context, C_exit_code);
+}

--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/src/CodeGear.h	Thu Jan 16 15:12:06 2020 +0900
@@ -0,0 +1,8 @@
+typedef struct CodeGear<Impl>{
+        union Data* codeGear;
+        enum Code code;
+        __code code(struct Integer* input1, struct Integer* input2, __code next(struct Integer* output, ...));
+        __code setInfo(struct Context* codeGear, union Data** dataGears, __code next(...));
+        union Data* dataGears[10];
+        __code next(...);
+} CodeGear;

--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/src/Executor.h	Thu Jan 16 15:12:06 2020 +0900
@@ -0,0 +1,8 @@
+typedef struct Executor<Impl>{
+    union Data* Executor;
+    struct Context* task;
+    __code next(...);
+    __code read(Impl* executor, struct Context* task, __code next(...));
+    __code exec(Impl* executor, struct Context* task, __code next(...));
+    __code write(Impl* executor, struct Context* task, __code next(...));
+} Executor;

--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/src/Iterator.h	Thu Jan 16 15:12:06 2020 +0900
@@ -0,0 +1,9 @@
+typedef struct Iterator<Impl>{
+        union Data* iterator;
+        struct Context* task;
+        int numGPU;
+        __code exec(Impl* iterator, struct Context* task, int numGPU, __code next(...));
+        __code barrier(Impl* iterator, struct Context* task, __code next(...), __code whenWait(...));
+        __code whenWait(...);
+        __code next(...);
+} Iterator;

--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/src/Lock.h	Thu Jan 16 15:12:06 2020 +0900
@@ -0,0 +1,6 @@
+typedef struct Lock<Impl>{
+        union Data* lock;
+        __code doLock(Impl* lock, __code next(...)); 
+        __code doUnlock(Impl* lock, __code next(...)); 
+        __code next(...);
+} Lock;

--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/src/LockImpl.cbc	Thu Jan 16 15:12:06 2020 +0900
@@ -0,0 +1,86 @@
+#include "../context.h"
+#interface "Queue.h"
+#interface "Atomic.h"
+#interface "Lock.h"
+#interface "Worker.h"
+#interface "TaskManager.h"
+
+Lock* createLockImpl(struct Context* context) {
+    struct Lock* lock = new Lock();
+    struct LockImpl* lockImpl = new LockImpl();
+    lockImpl->lock = NULL;
+    lockImpl->waitThreadQueue = createSynchronizedQueue(context);
+    lockImpl->atomic = createAtomicReference(context);
+    lock->lock = (union Data*)lockImpl;
+    lock->doLock = C_doLockLockImpl;
+    lock->doUnlock = C_doUnlockLockImpl;
+    return lock;
+}
+
+__code doLockLockImpl(struct LockImpl* lock, __code next(...)) {
+    struct Atomic* atomic = lock->atomic;
+    goto atomic->checkAndSet(&lock->lock, NULL, 1, doLockLockImpl1, doLockLockImpl2);
+}
+
+__code doLockLockImpl1(struct LockImpl* lock, __code next(...)) {
+    lock->lockContext = context;
+    goto next(...);
+}
+
+__code doLockLockImpl2(struct LockImpl* lock, __code next(...)) {
+    struct Queue* queue = lock->waitThreadQueue;
+    context->next= C_doLockLockImpl;
+    printf("Put task\n");
+    goto queue->put(context, doLockLockImpl3);
+}
+
+__code doLockLockImpl3(struct LockImpl* lock, struct Worker* worker, __code next(...)) {
+    goto worker->taskReceive(); // goto shceduler
+}
+
+__code doLockLockImpl3_stub(struct Context* context) {
+    // switch worker context
+    struct Context* workerContext = context->worker->worker->CPUWorker.context;
+    LockImpl* lockImpl = (LockImpl*)GearImpl(context, Lock, lock);
+    goto doLockLockImpl3(workerContext,
+            lockImpl,
+            context->worker,
+            Gearef(context, Lock)->next);
+}
+
+__code doUnlockLockImpl(struct LockImpl* lock, __code next(...)) {
+    if (lock->lockContext == context) {
+        struct Atomic* atomic = lock->atomic;
+        goto atomic->checkAndSet(&lock->lock, 1, NULL, doUnlockLockImpl1, doUnlockLockImpl);
+    }
+    goto next(...);
+}
+
+__code doUnlockLockImpl1(struct LockImpl* lock, __code next(...)) {
+    struct Queue* queue = lock->waitThreadQueue;
+    goto queue->isEmpty(doUnlockLockImpl2, doUnlockLockImpl4);
+}
+
+__code doUnlockLockImpl2(struct LockImpl* lock, __code next(...)) {
+    struct Queue* queue = lock->waitThreadQueue;
+    printf("%p: Take task\n", lock);
+    goto queue->take(doUnlockLockImpl3);
+}
+
+__code doUnlockLockImpl3(struct LockImpl* lock, struct Context* waitTask, __code next(...)) {
+    struct TaskManager* taskManager = waitTask->taskManager;
+    goto taskManager->spawn(waitTask, next(...)); //notify
+}
+
+__code doUnlockLockImpl3_stub(struct Context* context) {
+    LockImpl* lockImpl = (LockImpl*)GearImpl(context, Lock, lock);
+    struct Context* waitTask = &Gearef(context, Queue)->data->Context;
+    goto doUnlockLockImpl3(context,
+            lockImpl,
+            waitTask,
+            Gearef(context, Lock)->next);
+}
+
+__code doUnlockLockImpl4(struct LockImpl* lock, __code next(...)) {
+    goto next(...);
+}

--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/src/MultiDimIterator.cbc	Thu Jan 16 15:12:06 2020 +0900
@@ -0,0 +1,96 @@
+#include "../context.h"
+#interface "Iterator.h"
+#interface "TaskManager.h"
+#include <stdio.h>
+
+Iterator* createMultiDimIterator(struct Context* context, int x, int y, int z) {
+    struct Iterator* iterator = new Iterator();
+    struct MultiDimIterator* multiDimIterator = new MultiDimIterator();
+    iterator->iterator = (union Data*)multiDimIterator;
+    iterator->exec = C_execMultiDimIterator;
+    iterator->barrier = C_barrierMultiDimIterator;
+    multiDimIterator->x = x;
+    multiDimIterator->y = y;
+    multiDimIterator->z = z;
+    multiDimIterator->count = x * y * z;
+    multiDimIterator->counterX = 0;
+    multiDimIterator->counterY = 0;
+    multiDimIterator->counterZ = 0;
+    return iterator;
+}
+
+/**
+ * create iterateTask with index, that copy from task argument
+ * @return created iterateTask
+ * @param task      task of the copy source
+ * @x     index
+ */
+struct Context* createMultiDimIterateTask(struct Context* task, int x, int y, int z) {
+    struct Context* task1 = NEW(struct Context);
+    initContext(task1);
+    task1->taskManager = task->taskManager;
+    task1->next     = task->next;
+    task1->iterate  = 1;
+    task1->iterator = task->iterator;
+    task1->idgCount = task->idgCount;
+    task1->idg      = task->idg;
+    task1->maxIdg   = task->maxIdg;
+    for(int i = task1->idg; i < task1->maxIdg; i++) {
+        task1->data[i] = task->data[i];
+    }
+
+    // create index data gear and register input data to iterate task
+    struct MultiDim* multiDim = &ALLOCATE_DATA_GEAR(task1, MultiDim)->MultiDim;
+    multiDim->x = x;
+    multiDim->y = y;
+    multiDim->z = z;
+    task1->data[task1->maxIdg++] = (union Data*)multiDim;
+    task1->odg      = task->odg + 1;
+    task1->maxOdg   = task->maxOdg + 1;
+    for (int i = task1->odg; i < task1->maxOdg; i++) {
+        task1->data[i] = task->data[i-1];
+    }
+
+    return task1;
+}
+
+__code execMultiDimIterator(struct MultiDimIterator* iterator, struct Context* task, int numGPU, __code next(...)) {
+    // No GPU device
+    if (numGPU == 0) {
+        goto execMultiDimIterator1();
+    }
+    task->iterate = 1;
+    task->gpu = 1;
+    struct TaskManager* taskManager = task->taskManager;
+    goto taskManager->spawn(task, next(...));
+}
+
+__code execMultiDimIterator1(struct MultiDimIterator* iterator, struct Context* task, __code next(...)) {
+    int x = iterator->counterX;
+    int y = iterator->counterY;
+    int z = iterator->counterZ;
+    struct Context* iterateTask = createMultiDimIterateTask(task, x, y, z);
+    struct TaskManager* taskManager = task->taskManager;
+    goto taskManager->spawn(iterateTask, execMultiDimIterator2);
+}
+
+__code execMultiDimIterator2(struct MultiDimIterator* iterator, struct Context* task, __code next(...)) {
+    if (++iterator->counterX >= iterator->x) {
+        iterator->counterX = 0;
+        if (++iterator->counterY >= iterator->y) {
+            iterator->counterY = 0;
+            if (++iterator->counterZ >= iterator->z) {
+                iterator->counterZ = 0;
+                goto next(...);
+            }
+        }
+    }
+    goto execMultiDimIterator1();
+}
+
+__code barrierMultiDimIterator(struct MultiDimIterator* iterator, struct Context* task, __code next(...), __code whenWait(...)) {
+    if (task->gpu || __sync_fetch_and_sub(&iterator->count, 1) == 1) {
+        goto next(...);
+    }
+    goto whenWait(...);
+}

--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/src/Queue.h	Thu Jan 16 15:12:06 2020 +0900
@@ -0,0 +1,10 @@
+typedef struct Queue<Impl>{
+        union Data* queue;
+        union Data* data;
+        __code whenEmpty(...);
+        __code clear(Impl* queue, __code next(...));
+        __code put(Impl* queue, union Data* data, __code next(...));
+        __code take(Impl* queue, __code next(union Data*, ...));
+        __code isEmpty(Impl* queue, __code next(...), __code whenEmpty(...));
+        __code next(...);
+} Queue;

--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/src/RedBlackTree.agda	Thu Jan 16 15:12:06 2020 +0900
@@ -0,0 +1,145 @@
+module RedBlackTree where
+
+open import stack
+open import Level
+
+record TreeMethods {n m : Level } {a : Set n } {t : Set m } (treeImpl : Set n ) : Set (m Level.⊔ n) where
+  field
+    putImpl : treeImpl -> a -> (treeImpl -> t) -> t
+    getImpl  : treeImpl -> (treeImpl -> Maybe a -> t) -> t
+open TreeMethods
+
+record Tree  {n m : Level } {a : Set n } {t : Set m } (treeImpl : Set n ) : Set (m Level.⊔ n) where
+  field
+    tree : treeImpl
+    treeMethods : TreeMethods {n} {m} {a} {t} treeImpl
+  putTree : a -> (Tree treeImpl -> t) -> t
+  putTree d next = putImpl (treeMethods ) tree d (\t1 -> next (record {tree = t1 ; treeMethods = treeMethods} ))
+  getTree : (Tree treeImpl -> Maybe a -> t) -> t
+  getTree next = getImpl (treeMethods ) tree (\t1 d -> next (record {tree = t1 ; treeMethods = treeMethods} ) d )
+
+open Tree
+
+data Color {n : Level } : Set n where
+  Red   : Color
+  Black : Color
+
+data CompareResult {n : Level } : Set n where
+  LT : CompareResult
+  GT : CompareResult
+  EQ : CompareResult
+
+record Node {n : Level } (a k : Set n) : Set n where
+  inductive
+  field
+    key   : k
+    value : a
+    right : Maybe (Node a k)
+    left  : Maybe (Node a k)
+    color : Color {n}
+open Node
+
+record RedBlackTree {n m : Level } {t : Set m} (a k si : Set n) : Set (m Level.⊔ n) where
+  field
+    root : Maybe (Node a k)
+    nodeStack : Stack {n} {m} (Node a k) {t} si
+    compare : k -> k -> CompareResult {n}
+
+open RedBlackTree
+
+open Stack
+
+--
+-- put new node at parent node, and rebuild tree to the top
+--
+{-# TERMINATING #-}   -- https://agda.readthedocs.io/en/v2.5.3/language/termination-checking.html
+replaceNode : {n m : Level } {t : Set m } {a k si : Set n} -> RedBlackTree {n} {m} {t} a k si -> Stack (Node a k) si -> Node a k -> Node a k -> (RedBlackTree {n} {m} {t} a k si -> t) -> t
+replaceNode {n} {m} {t} {a} {k} {si} tree s parent n0 next = popStack s (
+      \s grandParent -> replaceNode1 s grandParent ( compare tree (key parent) (key n0) ) )
+  where
+        replaceNode1 : Stack (Node a k) si -> Maybe ( Node a k ) -> CompareResult -> t
+        replaceNode1 s Nothing LT = next ( record tree { root = Just ( record parent { left = Just n0 ; color = Black } ) } )   
+        replaceNode1 s Nothing GT = next ( record tree { root = Just ( record parent { right = Just n0 ; color = Black } ) } )   
+        replaceNode1 s Nothing EQ = next ( record tree { root = Just ( record parent { right = Just n0 ; color = Black } ) } )   
+        replaceNode1 s (Just grandParent) result with result
+        ... | LT =  replaceNode tree s grandParent ( record parent { left = Just n0 } ) next
+        ... | GT =  replaceNode tree s grandParent ( record parent { right = Just n0 } ) next
+        ... | EQ =  next tree 
+
+rotateRight : {n m : Level } {t : Set m } {a k si : Set n} -> RedBlackTree {n} {m} {t} a k si -> Stack (Node  a k) {t} si -> Node a k -> Node a k -> Node a k -> (RedBlackTree {n} {m} {t} a k si -> t) -> t
+rotateRight {n} {m} {t} {a} {k} {si} tree s n0 parent grandParent next = {!!}
+
+rotateLeft : {n m : Level } {t : Set m } {a k si : Set n} -> RedBlackTree {n} {m} {t} a k si -> Stack (Node a k) {t} si -> Node a k -> Node a k -> Node a k -> (RedBlackTree {n} {m} {t} a k si -> t) -> t
+rotateLeft {n} {m} {t} {a} {k} {si} tree s n0 parent grandParent next = {!!}
+
+insertCase5 : {n m : Level } {t : Set m } {a k si : Set n} -> RedBlackTree {n} {m} {t} a k si -> Stack (Node a k) {t} si -> Node a k -> Node a k -> Node a k -> (RedBlackTree {n} {m} {t} a k si -> t) -> t
+insertCase5 {n} {m} {t} {a} {k} {si} tree s n0 parent grandParent next = {!!}
+
+insertCase4 : {n m : Level } {t : Set m } {a k si : Set n} -> RedBlackTree {n} {m} {t} a k si -> Stack (Node a k) {t} si -> Node a k -> Node a k -> Node a k -> (RedBlackTree {n} {m} {t} a k si -> t) -> t
+insertCase4 {n} {m} {t} {a} {k} {si} tree s n0 parent grandParent next = {!!}
+
+{-# TERMINATING #-}
+insertNode : {n m : Level } {t : Set m } {a k si : Set n} -> RedBlackTree {n} {m} {t} a k si -> Stack (Node a k) {t} si -> Node a k -> (RedBlackTree {n} {m} {t} a k si -> t) -> t
+insertNode {n} {m} {t} {a} {k} {si} tree s n0 next = get2Stack s (\ s d1 d2 -> insertCase1 s n0 d1 d2 )
+   where
+    insertCase1 : Stack (Node a k) si -> Node a k -> Maybe (Node a k) -> Maybe (Node a k) -> t    -- placed here to allow mutual recursion
+          -- http://agda.readthedocs.io/en/v2.5.2/language/mutual-recursion.html
+    insertCase3 : Stack (Node a k) si -> Node a k -> Node a k -> Node a k -> t
+    insertCase3 s n0 parent grandParent with left grandParent | right grandParent
+    ... | Nothing | Nothing = insertCase4 tree s n0 parent grandParent next
+    ... | Nothing | Just uncle  = insertCase4 tree s n0 parent grandParent next
+    ... | Just uncle | _  with compare tree ( key uncle ) ( key parent )
+    ...                   | EQ =  insertCase4 tree s n0 parent grandParent next
+    ...                   | _ with color uncle
+    ...                           | Red = pop2Stack s ( \s p0 p1 -> insertCase1 s ( 
+           record grandParent { color = Red ; left = Just ( record parent { color = Black ; left = Just n0 } )  ; right = Just ( record uncle { color = Black } ) }) p0 p1 )
+    ...                           | Black = insertCase4 tree s n0 parent grandParent next
+    insertCase2 : Stack (Node a k) si -> Node a k -> Node a k -> Node a k -> t
+    insertCase2 s n0 parent grandParent with color parent
+    ... | Black = replaceNode tree s grandParent n0 next
+    ... | Red = insertCase3 s n0 parent grandParent
+    insertCase1 s n0 Nothing Nothing = next tree
+    insertCase1 s n0 Nothing (Just grandParent) = replaceNode tree s grandParent n0 next
+    insertCase1 s n0 (Just grandParent) Nothing = replaceNode tree s grandParent n0 next
+    insertCase1 s n0 (Just parent) (Just grandParent) = insertCase2 s n0 parent grandParent
+      where
+
+findNode : {n m : Level } {a k si : Set n} {t : Set m} -> RedBlackTree {n} {m} {t} a k si -> Stack (Node a k) si -> (Node a k) -> (Node a k) -> (RedBlackTree {n} {m} {t} a k si -> Stack (Node a k) si -> Node a k -> t) -> t
+findNode {n} {m} {a} {k} {si} {t} tree s n0 n1 next = pushStack s n1 (\ s -> findNode1 s n1)
+  where
+    findNode2 : Stack (Node a k) si -> (Maybe (Node a k)) -> t
+    findNode2 s Nothing = next tree s n0
+    findNode2 s (Just n) = findNode tree s n0 n next
+    findNode1 : Stack (Node a k) si -> (Node a k)  -> t
+    findNode1 s n1 with (compare tree (key n0) (key n1))
+    ...                                | EQ = next tree s n0 
+    ...                                | GT = findNode2 s (right n1)
+    ...                                | LT = findNode2 s (left n1)
+
+
+leafNode : {n : Level } {a k : Set n}  -> k -> a -> Node a k
+leafNode k1 value = record {
+    key   = k1 ;
+    value = value ;
+    right = Nothing ;
+    left  = Nothing ;
+    color = Black 
+  }
+
+putRedBlackTree : {n m : Level } {a k si : Set n} {t : Set m} -> RedBlackTree {n} {m} {t} a k si -> k -> a -> (RedBlackTree {n} {m} {t} a k si -> t) -> t
+putRedBlackTree {n} {m} {a} {k} {si} {t} tree k1 value next with (root tree)
+...                                | Nothing = next (record tree {root = Just (leafNode k1 value) })
+...                                | Just n2  = findNode tree (nodeStack tree) (leafNode k1 value) n2 (\ tree1 s n1 -> insertNode tree1 s n1 next)
+
+getRedBlackTree : {n m : Level } {a k si : Set n} {t : Set m} -> RedBlackTree {n} {m} {t} a k si -> k -> (RedBlackTree {n} {m} {t} a k si -> (Maybe (Node a k)) -> t) -> t
+getRedBlackTree {_} {_} {a} {k} {_} {t} tree k1 cs = checkNode (root tree)
+  where
+    checkNode : Maybe (Node a k) -> t
+    checkNode Nothing = cs tree Nothing
+    checkNode (Just n) = search n
+      where
+        search : Node a k -> t
+        search n with compare tree k1 (key n)
+        search n | LT = checkNode (left n)
+        search n | GT = checkNode (right n)
+        search n | EQ = cs tree (Just n)

--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/src/RedBlackTree.cbc	Thu Jan 16 15:12:06 2020 +0900
@@ -0,0 +1,602 @@
+#include <stdio.h>
+
+#include "../context.h"
+#interface "Tree.h"
+#interface "Stack.h"
+
+extern enum Relational compare(struct Node* node1, struct Node* node2);
+
+Tree* createRedBlackTree(struct Context* context) {
+    struct Tree* tree = new Tree();
+    struct RedBlackTree* redBlackTree = new RedBlackTree();
+    tree->tree = (union Data*)redBlackTree;
+    redBlackTree->root = NULL;
+    redBlackTree->nodeStack = createSingleLinkedStack(context);
+    tree->put = C_putRedBlackTree;
+    tree->get = C_getRedBlackTree;
+    tree->remove = C_removeRedBlackTree;
+    // tree->clear = C_clearRedBlackTree;
+    return tree;
+}
+
+void printTree1(union Data* data) {
+    struct Node* node = &data->Node;
+    if (node == NULL) {
+        printf("NULL");
+    } else {
+        printf("key = %d (", node->key);
+        printTree1((union Data*)(node->right));
+        printf("), (");
+        printTree1((union Data*)(node->left));
+        printf(")");
+    }
+}
+
+void printTree(union Data* data) {
+    printTree1(data);
+    printf("\n");
+}
+
+__code putRedBlackTree(struct RedBlackTree* tree, struct Node* node) {
+    struct Node* newNode = &ALLOCATE(context, Node)->Node;
+    struct Node* root = tree->root;
+    printTree((union Data*)(tree->root));
+    tree->newNode = newNode;
+    tree->root = newNode; // this should done at stackClear
+    tree->parent = NULL;
+    if (root) {
+        tree->current = root;
+        tree->result = compare(tree->current, node);
+        tree->findNodeNext = C_insertNode;
+        goto findNode(tree);
+    }
+    goto insertNode(tree, node);
+}
+
+__code findNode(struct RedBlackTree* tree) {
+    struct Stack* nodeStack = tree->nodeStack;
+    struct Node* oldNode = tree->current;
+    struct Node* newNode = tree->newNode;
+    tree->previous = newNode;
+    *newNode = *oldNode;
+    goto nodeStack->push((union Data*)newNode, findNode1);
+}
+
+__code findNode1(struct RedBlackTree* tree, struct Node* node, __code next(...)) {
+    struct Node* oldNode = tree->current;
+    struct Node* newNode = tree->previous;
+    struct Node* newnewNode = &ALLOCATE(context, Node)->Node;
+    int result = tree->result;
+    if (result == EQ) {
+        newNode->value = node->value;
+        // go to stack clear
+        goto next(...);
+    } else if (result == GT) {
+        tree->current = oldNode->right;
+        newNode->right = newnewNode;
+    } else {
+        tree->current = oldNode->left;
+        newNode->left = newnewNode;
+    }
+    tree->newNode = newnewNode;
+    if (tree->current) {
+        tree->result = compare(tree->current, node);
+        goto findNode(tree);
+    }
+    goto meta(context, tree->findNodeNext);
+    //   gato tree->findNodeNext(tree, node);
+    
+}
+
+__code insertNode(struct RedBlackTree* tree, struct Node* node) {
+    struct Stack* nodeStack = tree->nodeStack;
+    struct Node* newNode = tree->newNode;
+    *newNode = *node;
+    newNode->color = Red;
+    tree->current = newNode;
+    goto nodeStack->get2(insertCase1);
+}
+
+__code insertCase1(struct RedBlackTree* tree, struct Node *parent, struct Node *grandparent) {
+    if (parent != NULL) {
+        tree->parent = parent;
+        tree->grandparent = grandparent;
+        goto insertCase2(tree);
+    }
+    tree->root->color = Black;
+    goto stackClear();
+}
+
+__code insertCase1_stub(struct Context* context) {
+    goto insertCase1(context, 
+        &Gearef(context, Tree)->tree->Tree.tree->RedBlackTree,
+        &context->data[D_Stack]->Stack.data->Node,
+        &context->data[D_Stack]->Stack.data1->Node);
+}
+
+__code insertCase2(struct RedBlackTree* tree) {
+    if (tree->parent->color == Black) {
+        goto stackClear();
+    }
+    goto insertCase3(tree);
+}
+
+__code insertCase3(struct RedBlackTree* tree) {
+    struct Stack* nodeStack = tree->nodeStack;
+    struct Node* uncle;
+
+    if (tree->grandparent->left == tree->parent) {
+        uncle = tree->grandparent->right;
+    } else {
+        uncle = tree->grandparent->left;
+    }
+
+    if (uncle && (uncle->color == Red)) {
+        // do insertcase1 on grandparent, stack must be pop by two
+        tree->parent->color = Black;
+        uncle->color = Black;
+        tree->grandparent->color = Red;
+        tree->current = tree->grandparent;
+        goto nodeStack->pop2(insertCase1);
+    }
+    goto insertCase4();
+}
+
+__code insertCase4(struct RedBlackTree* tree, struct RotateTree* rotateTree) {
+    struct Stack* nodeStack = tree->nodeStack;
+
+    if ((tree->current == tree->parent->right) && (tree->parent == tree->grandparent->left)) {
+        tree->current = tree->current->left;
+        tree->parent = tree->grandparent;
+
+        rotateTree->traverse = tree;
+        rotateTree->next = C_insertCase5;
+
+        goto nodeStack->pop(rotateLeft);
+    } else if ((tree->current == tree->parent->left) && (tree->parent == tree->grandparent->right)) {
+        tree->parent = tree->grandparent;
+        tree->current = tree->current->right;
+
+        rotateTree->traverse = tree;
+        rotateTree->next = C_insertCase5;
+
+        goto nodeStack->pop(rotateRight);
+    }
+
+    goto insertCase5();
+}
+
+__code insertCase5(struct RedBlackTree* tree) {
+    struct Stack* nodeStack = tree->nodeStack;
+    goto nodeStack->pop2(insertCase51);
+}
+
+__code insertCase51(struct RedBlackTree* tree, struct RotateTree* rotateTree, struct Node* parent, struct Node* grandparent) {
+    struct Node* current = tree->current;
+    tree->parent = parent;
+    tree->grandparent = grandparent;
+
+    parent->color = Black;
+    grandparent->color = Red;
+
+    tree->current = grandparent;
+
+    rotateTree->traverse = tree;
+    rotateTree->next = C_stackClear;
+
+    if ((current == parent->left) && (parent == grandparent->left)){
+        goto rotateRight();
+    } else {
+        goto rotateLeft();
+    }
+}
+
+__code insertCase51_stub(struct Context* context) {
+    struct Node* parent = &context->data[D_Stack]->Stack.data->Node;
+    struct Node* grandparent = &context->data[D_Stack]->Stack.data1->Node;
+    goto insertCase51(context,
+                      &Gearef(context, Tree)->tree->Tree.tree->RedBlackTree,
+                      Gearef(context, RotateTree),
+                      parent,
+                      grandparent);
+}
+
+__code rotateLeft(struct RedBlackTree* tree) {
+    struct Stack* nodeStack = tree->nodeStack;
+    goto nodeStack->get(rotateLeft1);
+}
+
+__code rotateLeft_stub(struct Context* context) {
+    struct RedBlackTree* traverse = context->data[D_RotateTree]->RotateTree.traverse;
+    goto rotateLeft(context, traverse);
+}
+    
+__code rotateLeft1(struct Node* node, struct RedBlackTree* tree, struct Node* parent, struct RotateTree* rotateTree) {
+    struct Node* tmp = node->right;
+
+    if (parent) {
+        if (node == parent->left)
+            parent->left = tmp;
+        else
+            parent->right = tmp;
+    } else {
+        tree->root = tmp;
+    }
+
+    node->right = tmp->left;
+    tmp->left = node;
+    tree->current = tmp;
+    
+    goto meta(context, rotateTree->next);
+}
+
+__code rotateLeft1_stub(struct Context* context) {
+    struct RedBlackTree* traverse = context->data[D_RotateTree]->RotateTree.traverse;
+    struct Node* parent = &context->data[D_Stack]->Stack.data->Node;
+    goto rotateLeft1(context,
+                    traverse->current,
+                    traverse,
+                    parent,
+                    Gearef(context, RotateTree));
+}
+
+__code rotateRight(struct RedBlackTree* tree) {
+    struct Stack* nodeStack = tree->nodeStack;
+    goto nodeStack->get(rotateRight1);
+}
+
+__code rotateRight_stub(struct Context* context) {
+    struct RedBlackTree* traverse = context->data[D_RotateTree]->RotateTree.traverse;
+    goto rotateLeft(context, traverse);
+}
+
+__code rotateRight1(struct Node* node, struct RedBlackTree* traverse,struct Node *parent,struct RotateTree *rotateTree) {
+    struct Node* tmp = node->left;
+    
+    if (parent) {
+        if (node == parent->left)
+            parent->left = tmp;
+        else
+            parent->right = tmp;
+    } else {
+        traverse->root = tmp;
+    }
+
+    node->left = tmp->right;
+    tmp->right = node;
+    traverse->current = tmp;
+    
+    goto meta(context, rotateTree->next);
+}
+
+__code rotateRight1_stub(struct Context* context) {
+    struct RedBlackTree* traverse = context->data[D_RotateTree]->RotateTree.traverse;
+    struct Node* parent = &context->data[D_Stack]->Stack.data->Node;
+    goto rotateRight1(context,
+                     traverse->current,
+                     traverse,
+                     parent,
+                     Gearef(context, RotateTree));
+}
+
+__code stackClear(struct RedBlackTree* tree, struct Stack* nodeStack, __code next(...)) {
+    tree->current = 0;
+    nodeStack->stack = (union Data*)tree->nodeStack;
+    nodeStack->next = next;
+    goto meta(context, tree->nodeStack->clear);
+}
+
+__code getRedBlackTree(struct RedBlackTree* tree, __code next(...)) {
+    if (tree->root) {
+        tree->current = tree->root;
+
+        goto search();
+    }
+
+    goto next(...);
+}
+
+__code search(struct RedBlackTree* tree, struct Node* node, __code next(...)) {
+    // compare(context, traverse, traverse->current->key, node->key);
+    tree->result = compare(tree->current, node);
+    if (tree->result == EQ) {
+        *node = *tree->current;
+        
+        goto meta(context, next);
+    } else if (tree->result == GT) {
+        tree->current = tree->current->right;
+    } else {
+        tree->current = tree->current->left;
+    }
+        
+    if (tree->current) {
+        goto meta(context, C_search);
+    }
+
+    goto next(...);
+}
+
+
+__code removeRedBlackTree(struct RedBlackTree* tree, struct Node* node, __code next(...)) {
+    struct Node* newNode = &ALLOCATE(context, Node)->Node;
+    struct Node* root = tree->root;
+    printTree((union Data*)(tree->root));
+    tree->newNode = newNode;
+    tree->root = newNode; // this should done at stackClear
+    tree->parent = NULL;
+    if (root) {
+        tree->current = root;
+        tree->result = compare(tree->current, node);
+        tree->findNodeNext = C_replaceNodeForDelete2;
+        goto findNode(tree);
+    }
+    goto next(...);
+}
+
+
+
+__code delete2(struct Node* current) {
+    if (current->color == Black) {
+        struct Node* child = current->right == NULL ? current->left : current->right;
+        current->color = child == NULL ? Black : child->color;
+
+        goto deleteCase1(current);
+    }
+
+    goto delete3(tree, current);
+}
+
+
+
+__code delete3(struct RedBlackTree* tree, struct Node* current, __code next(...)) {
+    struct Node* tmp = current->right == NULL ? current->left : current->right;
+    struct Stack* nodeStack = tree->nodeStack;
+
+    if (tree->parent) {
+        if (current == tree->parent->left)
+            tree->parent->left = tmp;
+        else
+            tree->parent->right = tmp;
+    } else {
+        tree->root = tmp;
+    }
+
+
+    if (tree->parent == NULL && tmp) {
+        tmp->color = Black;
+    }
+
+    current == tree->parent->left ? (tree->parent->left = NULL) : (tree->parent->right = NULL);
+
+    Gearef(context, Stack)->stack = (union Data*) nodeStack;
+    Gearef(context, Stack)->next = next;
+    goto meta(context, nodeStack->pop);
+
+//    gato nodeStack->pop(next);
+}
+
+
+
+__code replaceNodeForDelete2(struct RedBlackTree* tree, struct Node* newNode) {
+    if (tree->current->left && tree->current->right) {
+        tree->parent = newNode;
+        tree->current = newNode->left;
+        newNode->left = context->heap;
+
+
+        tree->parent = newNode;
+        
+        goto findMax1(tree,oldNode, newNode);
+    }
+
+    goto delete2(current);
+}
+
+
+__code findMax1(struct RedBlackTree* tree, struct Node* oldNode, struct Node* newNode) {
+    *newNode = *oldNode;
+
+    if (newNode->right) {
+        goto findMax2(tree, oldNode, newNode);
+    }
+    
+    tree->current = newNode;
+
+    goto delete2(current);
+}
+
+
+    
+
+__code findMax2(struct RedBlackTree* tree, struct Node* oldNode, struct Node* newNode) {
+    *newNode = *oldNode;
+
+    if (newNode->right->right) {
+        tree->current = newNode->right;
+        newNode->right = context->heap;
+
+        tree->parent = newNode;
+        
+        goto findMax2(tree, oldNode, newNode);
+    }
+
+    tree->current = newNode;
+    
+    goto delete2(tree,current);
+}
+    
+
+__code deleteCase1(struct RedBlackTree* tree, struct Node* current) {
+    if (tree->parent) {
+        goto deleteCase2(tree,current);
+    }
+
+    goto delete3(tree, current);
+}
+
+
+
+__code deleteCase2(struct RedBlackTree* tree, struct Node* current, struct RotateTree* rotateTree) {
+    struct Node* sibling = current == tree->parent->left ? tree->parent->right : tree->parent->left;
+    struct Stack* nodeStack = tree->nodeStack;
+    
+    if ((sibling == NULL ? Black : sibling->color) == Red) {
+        tree->parent->color = Red;
+        sibling->color = Black;
+
+        current == tree->parent->left ? (tree->parent->left = context->heap) : (tree->parent->right = context->heap);
+
+        struct Node* node = sibling;
+        
+        tree->current = tree->parent;
+
+        rotateTree->traverse = tree;
+        rotateTree->next = C_deleteCase3;
+
+        if (current == tree->parent->left) {
+            goto nodeStack->push((union Data*)node,rotateLeft);
+        } else {
+            goto nodeStack->push((union Data*)node,rotateRight);
+        }
+
+        goto deleteCase3(tree,current);
+    }
+}
+
+
+
+__code deleteCase3(struct RedBlackTree* tree, struct Node* current) {
+    struct Node* sibling = current == tree->parent->left ? tree->parent->right : tree->parent->left;
+    
+    if (tree->parent->color == Black &&
+        (sibling == NULL ? Black : sibling->color) == Black &&
+        (sibling->left == NULL ? Black : sibling->left->color) == Black &&
+        (sibling->right == NULL ? Black : sibling->right->color) == Black) {
+        sibling->color = Red;
+
+        tree->current = tree->parent;
+        goto deleteCase1(current);
+    }
+
+    goto deleteCase4(current);
+}
+
+
+
+__code deleteCase4(struct RedBlackTree* tree,struct Node* current) {
+    struct Node* sibling = current == tree->parent->left ? tree->parent->right : tree->parent->left;
+    
+    if (tree->parent->color == Red &&
+        (sibling == NULL ? Black : sibling->color) == Black &&
+        (sibling->left == NULL ? Black : sibling->left->color) == Black &&
+        (sibling->right == NULL ? Black : sibling->right->color) == Black) {
+        sibling->color = Red;
+        tree->parent->color = Black;
+
+        goto delete3(tree,current);
+    }
+
+    goto deleteCase5(tree,current);
+}
+
+
+
+__code deleteCase5(struct RedBlackTree* tree, struct Node* current, struct RotateTree* rotateTree) {
+    struct Node* sibling = current == tree->parent->left ? tree->parent->right : tree->parent->left;
+    struct Stack* nodeStack = tree->nodeStack;
+    // sibling->parent = tree->parent;
+    
+    if (current == tree->parent->left &&
+        (sibling == NULL ? Black : sibling->color) == Black &&
+        (sibling->left == NULL ? Black : sibling->left->color) == Red &&
+        (sibling->right == NULL ? Black : sibling->right->color) == Black) {
+        sibling->color = Red;
+        sibling->left->color = Black;
+        
+        // sibling == sibling->parent->left ? (sibling->parent->left = context->heap) : (sibling->parent->right = context->heap);
+        sibling == tree->parent->left ? (tree->parent->left = context->heap) : (tree->parent->right = context->heap);
+
+        struct Node* node = new Node();
+        node = sibling->left;
+
+        struct Node* tmp = node;
+        *tmp = *sibling;
+        tree->parent = current;
+        
+        tmp->left = context->heap;
+/*         struct Node* node = new Node(); */
+/*         node = *sibling->left; */
+        tree->parent = tmp;
+
+        tree->current = tmp;
+        
+
+        rotateTree->traverse = tree;
+        rotateTree->next = C_deleteCase6;
+
+        goto nodeStack->push((union Data*)node,rotateRight);
+    } else if (current == tree->parent->right &&
+               (sibling == NULL ? Black : sibling->color) == Black &&
+               (sibling->left == NULL ? Black : sibling->left->color) == Black &&
+               (sibling->right == NULL ? Black : sibling->right->color) == Red) {
+        sibling->color = Red;
+        sibling->right->color = Black;
+
+        sibling == tree->parent->left ? (tree->parent->left = context->heap) : (tree->parent->right = context->heap);
+
+        struct Node* node = new Node();
+        node = sibling->right;
+
+        struct Node* tmp = node;
+        *tmp = *sibling;
+        // tmp->parent = current;
+
+        tmp->right = context->heap;
+/*         struct Node* node = new Node(); */
+/*         node = *sibling->right; */
+        //node->parent = tmp;
+
+        tree->current = tmp;
+        
+
+        rotateTree->traverse = tree;
+        rotateTree->next = C_deleteCase6;
+
+        goto nodeStack->push((union Data*)node,rotateLeft);
+    }
+
+    goto deleteCase6(tree,current);
+}
+
+
+__code deleteCase6(struct RedBlackTree* tree, struct Node* current, struct RotateTree* rotateTree) {
+    struct Node* sibling = current == tree->parent->left ? tree->parent->right : tree->parent->left;
+    struct Stack* nodeStack = tree->nodeStack;
+    sibling == tree->parent->left ? (tree->parent->left = context->heap) : (tree->parent->right = context->heap);
+
+    struct Node* tmp = sibling;
+    // *tmp = *sibling;
+    tree->parent = current;
+
+    tmp->color = tree->parent->color;
+    tree->parent->color = Black;
+    
+    
+    if (current == tree->parent->left) {
+        tmp->right->color = Black;
+        tree->current = tree->parent;
+
+        rotateTree->traverse = tree;
+        rotateTree->next = C_delete3;
+
+        goto nodeStack->push((union Data*)tmp,rotateLeft);
+    } else {
+        tmp->left->color = Black;
+        tree->current = tree->parent;
+
+        rotateTree->traverse = tree;
+        rotateTree->next = C_delete3;
+
+        goto nodeStack->push((union Data*)tmp,rotateLeft);
+    }
+}

--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/src/RedBlackTreeReWright.cbc	Thu Jan 16 15:12:06 2020 +0900
@@ -0,0 +1,269 @@
+#include <stdio.h>
+
+#include "../context.h"
+#include "../compare.c"
+#interface "Tree.h"
+#interface "Stack.h"
+
+extern enum Relational compare(struct Node* node1, struct Node* node2);
+
+
+Tree* createRedBlackTree(struct Context* context) {
+    struct Tree* tree = new Tree();
+    struct RedBlackTree* rbtree = new RedBlackTree();
+
+    tree->tree = (union Data*)rbtree;
+     rbtree->root = NULL;
+     rbtree->nodeStack = (union Data*)createSingleLinkedStack(context);
+     tree->put = C_putRedBlackTree;
+    // tree->get = C_getRedBlackTree;
+    // tree->remove = C_removeRedBlackTree;
+    // tree->clear = C_clearRedBlackTree;
+     return tree;
+}
+
+void printNode(struct Node* node) {
+  if (node == NULL) {
+    printf("leaf");
+  } else {
+    printf("((%d,%d (",node->color, node->key);
+    printNode(node->right);
+    printf(") (");
+    printNode(node->left);
+    printf(")");
+  }
+}
+
+void printTree(struct RedBlackTree* tree) {
+  printf("\n");
+  tree->current = tree->root;
+  printNode(tree->current);
+  printf(")\n");
+}
+
+__code putRedBlackTree(struct RedBlackTree* tree, struct Node* node, __code next(...)) {
+    printf("C_putRedBlackTree\n");
+    printf("value->%d,key->%d \n",node->value,node->key);
+    tree->previous = tree->newNode;
+    tree->newNode = node;
+    tree->newNode->color = Red;
+    tree->current = tree->root;
+    goto insertRBTree(node, tree);
+}
+
+__code stackClear(struct RedBlackTree* tree, struct Stack* nodeStack, __code next(...)) {
+   tree->current = 0;
+   nodeStack->stack = tree->nodeStack;
+   nodeStack->next = next;
+   goto meta(context, tree->nodeStack->clear);
+  }
+
+__code getRedBlackTree(struct RedBlackTree* tree, __code next(...)) {
+    if (tree->root) {
+        tree->current = tree->root;
+        goto insertRBTree();
+        // goto deleteRBTree();
+      }
+    goto next(...);
+}
+
+__code insertRBTree(struct Node* node, struct RedBlackTree* tree, struct Stack* stack, __code next(...)) {
+  // first case tree->current = root;
+  printf("C_insertRBTree\n");
+  printf("value->%d,key->%d\n",node->value,node->key);
+  printf("newNode value->%d,newNode key->%d\n",tree->newNode->value,tree->newNode->key);
+
+  if (tree->root == NULL) {
+    printf("insertRBTree_root eq NULL\n");
+    tree->root = tree->newNode;
+    tree->root->color = Black;
+    printf("tree->root->key = %d, tree->root->color = %d \n",tree->root->key,tree->root->color);
+    printTree(tree);
+    goto next(tree,...);
+  } else {
+    goto searchInsertLocation(node, tree, stack);
+  }
+}
+
+__code insertRBTree_stub(struct Context* context) {
+	Node* node = Gearef(context, Tree)->node;
+	RedBlackTree* tree = (RedBlackTree*)GearImpl(context, Tree, tree);
+	Stack* stack = createSingleLinkedStack(context);
+	enum Code next = Gearef(context, Tree)->next;
+	goto insertRBTree(context, node, tree, stack, next);
+} 
+
+__code searchInsertLocation(struct Node* node, struct RedBlackTree* tree) {
+  // first case tree->current = root; PreCase remove root=NULL case.don't exist firstCase tree->current=NULL
+  printf("C_searchInsertLocation\n");
+  printf("nownode->key %d , previous->key %d \n",tree->newNode->key,tree->previous->key);
+
+  tree->result = compare(tree->current, node);
+  printf("tree->current->key = %d, node->key %d\n",tree->current->key,node->key);
+  printf("compare (%d,%d)\n",tree->current,node);
+
+  Stack* stack = tree->nodeStack;
+
+  if (tree->current == NULL) {
+    printf("goto insertLocationBackInsert stack->pop\n");
+    goto stack->pop(insertLocationBackInsert);
+  }
+  if (tree->result == GT) {
+    printf("GT searchInsertLocation\n");
+    tree->current = tree->current->right;
+    goto stack->push(tree->newNode,insertLocationBackInsert);
+  } else if (tree->result == LT) {
+    printf("LT searchInsertLocation\n");
+    tree->current = tree->current->left;
+    goto stack->push(tree->newNode, searchInsertLocation);
+  } else if (tree->result == EQ) {
+    printf("already member this node : __code searchInsertLocation()\n");
+    goto meta(context, C_exit_code);
+  } else {
+    printf("$insert value tree : __code searchInsertLocation() \n");
+    goto meta(context, C_exit_code);
+  }
+}
+
+__code searchInsertLocation_stub(struct Context* context) {
+	Node* node = Gearef(context, Tree)->node;
+	RedBlackTree* tree = (RedBlackTree*)GearImpl(context, Tree, tree);
+  Stack* stack = (struct Stack*)Gearef(context, Stack)->stack;
+	goto searchInsertLocation(context, node, tree);
+}
+
+__code insertLocationBackInsert(struct RedBlackTree* tree, struct Node* node, struct Stack* stack) {
+  printf("C_insertLocationBackInsert\n");
+  struct Node* hoge = stack->data;
+  printf("stackpopdata%d\n",stack->data);
+  tree->current = tree->previous;
+  // tree->current = nodeStack->data;
+  // this CS is ones only backTrace, and insert node
+  tree->result = compare(tree->previous,tree->newNode);
+  printf("back,compare\n");
+  if (tree->result == GT) {
+    printf("GT\n");
+    tree->current->right = tree->newNode;
+    printTree(tree);
+    goto insertBalance(tree, stack, node, next);
+  } else if (tree->result == LT) {
+    printf("LT\n");
+    tree->current->left = tree->newNode;
+    goto insertBalance(tree, stack, node, next);
+  } else {
+    printf("error : __code insertLocationBackTrace() \n");
+    goto meta(context, C_exit_code);
+  }
+}
+
+__code insertLocationBackInsert_stub(struct Context* context) {
+	RedBlackTree* tree = (RedBlackTree*)GearImpl(context, Tree, tree);
+  SingleLinkedStack* singleLinkedStack = (SingleLinkedStack*)GearImpl(context, Stack, stack);
+	Node* node = Gearef(context, Tree)->node;
+  Stack* stack = (struct Stack*)Gearef(context, Stack)->stack;
+	goto insertLocationBackInsert(context, tree, node, stack);
+}
+
+__code insertBalance(struct RedBlackTree* tree, struct Node* nodeStack, struct Node* node, __code next(...)) {
+  printf("C_insertBalance\n");
+  struct Node* traceNode = tree->nodeStack->data;
+  tree->current = traceNode;
+  struct Stack* stack = tree->nodeStack;
+
+  // exit insertion code
+  if (tree->current == tree->root) {
+    tree->current->color = Black;
+    printTree(tree);
+    //printTree
+    goto next(tree,...);
+  }
+
+
+  //current color eq Red
+  if (tree->current->color == Red)
+    goto stack->pop(insertBalance);
+
+  // current color eq Black
+  if (tree->current->left->left || tree->current->left->right) {
+    goto insertBalanceLeft(tree,nodeStack);
+  } else if (tree->current->right->left || tree->current->right->right) {
+    goto insertBalanceRight(tree,nodeStack);
+  } else {
+    goto stack->pop(insertBalance);
+  }
+}
+
+__code insertBalanceLeft(struct RedBlackTree* tree, struct Node* nodeStack, struct Node* node) {
+  printf("C_insertBalanceLeft\n");
+  struct Stack* stack = tree->nodeStack;
+
+  if (tree->current->color == Black && tree->current->left->color == Red && tree->current->left->left->color == Red) {
+    struct Node* tmpCurrent  = tree->current;
+    struct Node* tmpLeft     = tree->current->left;
+    struct Node* tmpLeftLeft = tree->current->left->left;
+
+    tree->current = tmpLeft;
+    tree->current->right = tmpCurrent;
+    tree->current->left = tmpLeftLeft;
+    tree->current->right->left = tmpLeft->right;
+    tree->current->color = Red;
+    tree->current->left->color = Black;
+    tree->current->right->color = Black;
+    goto stack->pop(insertBalance);
+
+  } else if (tree->current->color == Black && tree->current->left->color == Red && tree->current->left->right->color == Red) {
+    struct Node* tmpCurrent   = tree->current;
+    struct Node* tmpLeft      = tree->current->left;
+    struct Node* tmpLeftRight = tree->current->left->right;
+
+    tree->current = tmpLeft;
+    tree->current->right = tmpCurrent;
+    tree->current->left = tmpLeftRight;
+    tree->current->right->left = tmpLeft->left;
+    tree->current->color = Red;
+    tree->current->left->color = Black;
+    tree->current->right->color = Black;
+    goto stack->pop(insertBalance);
+
+  }
+}
+
+__code insertBalanceRight(struct RedBlackTree* tree, struct Node* nodeStack, struct Node* node) {
+  printf("C_insertBalanceLeft\n");
+  struct Stack* stack = tree->nodeStack;
+
+  if (tree->current->color == Black && tree->current->right->color == Red && tree->current->right->right->color == Red) {
+    struct Node* tmpCurrent    = tree->current;
+    struct Node* tmpRight      = tree->current->right;
+    struct Node* tmpRightRight = tree->current->right->right;
+
+    tree->current = tmpRight;
+    tree->current->left = tmpCurrent;
+    tree->current->right = tmpRightRight;
+    tree->current->left->right = tmpRight->left;
+    tree->current->color = Red;
+    tree->current->left->color = Black;
+    tree->current->right->color = Black;
+    goto stack->pop(insertBalance);
+
+  } else if (tree->current->color == Black && tree->current->right->color == Red && tree->current->right->left->color == Red) {
+
+    struct Node* tmpCurrent = tree->current;
+    struct Node* tmpRight = tree->current->right;
+    struct Node* tmpRightLeft = tree->current->right->left;
+
+    tree->current = tmpRight;
+    tree->current->right = tmpCurrent;
+    tree->current->left = tmpRightLeft;
+    tree->current->left->right = tmpRight->right;
+    tree->current->color = Red;
+    tree->current->left->color = Black;
+    tree->current->right->color = Black;
+    goto stack->pop(insertBalance);
+
+  } else {
+    printf("unkwon error : __code insertBalanceRight() \n");
+    goto meta(context, C_exit_code);
+  }
+}
+// insertCode end

--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/src/Semaphore.h	Thu Jan 16 15:12:06 2020 +0900
@@ -0,0 +1,6 @@
+typedef struct Semaphore<Impl>{
+        union Data* semaphore;
+        __code p(Impl* semaphore, __code next(...)); 
+        __code v(Impl* semaphore, __code next(...)); 
+        __code next(...);
+} Semaphore;

--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/src/SemaphoreImpl.cbc	Thu Jan 16 15:12:06 2020 +0900
@@ -0,0 +1,37 @@
+#include "../context.h"
+#interface "semaphore.h"
+
+Semaphore* createSemaphoreImpl(struct Context* context, int n) {
+    struct Semaphore* semaphore = new Semaphore();
+    struct SemaphoreImpl* semaphoreImpl = new SemaphoreImpl();
+    semaphore->semaphore = (union Data*)semaphoreImpl;
+    semaphoreImpl->value =  n;
+    pthread_mutex_init(&semaphoreImpl->mutex, NULL);
+    pthread_cond_init(&semaphoreImpl->cond, NULL);
+    semaphore->p = C_pOperationSemaphoreImpl;
+    semaphore->v = C_vOperationSemaphoreImpl;
+    return semaphore;
+}
+
+__code pOperationSemaphoreImpl(struct SemaphoreImpl* semaphore, __code next(...)) {
+    pthread_mutex_lock(&semaphore->mutex);
+    goto meta(context, C_pOperationSemaphoreImpl1);
+}
+
+__code pOperationSemaphoreImpl1(struct SemaphoreImpl* semaphore, __code next(...)) {
+    if (semaphore->value == 0) {
+        pthread_cond_wait(&semaphore->cond, &semaphore->mutex);
+        goto meta(context, C_pOperationSemaphoreImpl1);
+    }
+    semaphore->value--;
+    pthread_mutex_unlock(&semaphore->mutex);
+    goto next(...);
+}
+
+__code vOperationSemaphoreImpl(struct SemaphoreImpl* semaphore, __code next(...)) {
+    pthread_mutex_lock(&semaphore->mutex);
+    semaphore->value++;
+    pthread_cond_signal(&semaphore->cond);
+    pthread_mutex_unlock(&semaphore->mutex);
+    goto next(...);
+}

--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/src/SingleLinkedQueue.cbc	Thu Jan 16 15:12:06 2020 +0900
@@ -0,0 +1,68 @@
+#include "../context.h"
+#include <stdio.h>
+#interface "Queue.h"
+// use "Node.h"
+// use "Element.h"
+
+Queue* createSingleLinkedQueue(struct Context* context) {
+    struct Queue* queue = new Queue();
+    struct SingleLinkedQueue* singleLinkedQueue = new SingleLinkedQueue();
+    queue->queue = (union Data*)singleLinkedQueue;
+    singleLinkedQueue->top  = new Element();
+    singleLinkedQueue->last = singleLinkedQueue->top;
+    queue->take  = C_takeSingleLinkedQueue;
+    queue->put  = C_putSingleLinkedQueue;
+    queue->isEmpty = C_isEmptySingleLinkedQueue;
+    queue->clear = C_clearSingleLinkedQueue;
+    return queue;
+}
+
+void printQueue1(union Data* data) {
+    struct Node* node = &data->Element.data->Node;
+    if (node == NULL) {
+        printf("NULL");
+    } else {
+        printf("key = %d ,", node->key);
+        printQueue1((union Data*)data->Element.next);
+    }
+}
+
+void printQueue(union Data* data) {
+    printQueue1(data);
+    printf("\n");
+}
+
+__code clearSingleLinkedQueue(struct SingleLinkedQueue* queue, __code next(...)) {
+    queue->top = NULL;
+    goto next(...);
+}
+
+__code putSingleLinkedQueue(struct SingleLinkedQueue* queue, union Data* data, __code next(...)) {
+    Element* element = new Element();
+    element->data = data;
+    element->next = NULL;
+    queue->last->next  = element;
+    queue->last = element;
+    goto next(...);
+}
+
+__code takeSingleLinkedQueue(struct SingleLinkedQueue* queue, __code next(union Data* data, ...)) {
+    struct Element* top = queue->top;
+    struct Element* nextElement = top->next;
+    if (queue->top == queue->last) {
+        data = NULL;
+    } else {
+        queue->top = nextElement;
+        data = nextElement->data;
+    }
+    goto next(data, ...);
+}
+
+__code isEmptySingleLinkedQueue(struct SingleLinkedQueue* queue, __code next(...), __code whenEmpty(...)) {
+    if (queue->top == queue->last) {
+        goto whenEmpty(...);
+    } else {
+        goto next(...);
+    }
+}
+

--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/src/SingleLinkedStack.cbc	Thu Jan 16 15:12:06 2020 +0900
@@ -0,0 +1,111 @@
+#include "../context.h"
+#interface "Stack.h"
+#include <stdio.h>
+
+// typedef struct SingleLinkedStack {
+//     struct Element* top;
+// } SingleLinkedStack;
+
+Stack* createSingleLinkedStack(struct Context* context) {
+    struct Stack* stack = new Stack();
+    struct SingleLinkedStack* singleLinkedStack = new SingleLinkedStack();
+    stack->stack = (union Data*)singleLinkedStack;
+    singleLinkedStack->top = NULL;
+    stack->push = C_pushSingleLinkedStack;
+    stack->pop  = C_popSingleLinkedStack;
+    stack->pop2  = C_pop2SingleLinkedStack;
+    stack->get  = C_getSingleLinkedStack;
+    stack->get2  = C_get2SingleLinkedStack;
+    stack->isEmpty = C_isEmptySingleLinkedStack;
+    stack->clear = C_clearSingleLinkedStack;
+    return stack;
+}
+
+void printStack1(union Data* data) {
+    struct Node* node = &data->Element.data->Node;
+    if (node == NULL) {
+        printf("NULL");
+    } else {
+        printf("key = %d ,", node->key);
+        printStack1((union Data*)data->Element.next);
+    }
+}
+
+void printStack(union Data* data) {
+    printStack1(data);
+    printf("\n");
+}
+
+__code clearSingleLinkedStack(struct SingleLinkedStack* stack,__code next(...)) {
+    stack->top = NULL;
+    goto next(...);
+}
+
+__code pushSingleLinkedStack(struct SingleLinkedStack* stack, union Data* data, __code next(...)) {
+    Element* element = new Element();
+    element->next = stack->top;
+    element->data = data;
+    stack->top = element;
+    goto next(...);
+}
+
+__code popSingleLinkedStack(struct SingleLinkedStack* stack, __code next(union Data* data, ...)) {
+    if (stack->top) {
+        data = stack->top->data;
+        stack->top = stack->top->next;
+    } else {
+        data = NULL;
+    }
+    goto next(data, ...);
+}
+
+__code pop2SingleLinkedStack(struct SingleLinkedStack* stack, __code next(union Data* data, union Data* data1, ...)) {
+    if (stack->top) {
+        data = stack->top->data;
+        stack->top = stack->top->next;
+    } else {
+        data = NULL;
+    }
+    if (stack->top) {
+        data1 = stack->top->data;
+        stack->top = stack->top->next;
+    } else {
+        data1 = NULL;
+    }
+    goto next(data, data1, ...);
+}
+
+
+__code getSingleLinkedStack(struct SingleLinkedStack* stack, __code next(union Data* data, ...)) {
+    if (stack->top) {
+        data = stack->top->data;
+    } else {
+        data = NULL;
+    }
+    goto next(data, ...);
+}
+
+__code get2SingleLinkedStack(struct SingleLinkedStack* stack, __code next(union Data* data, union Data* data1, ...)) {
+    if (stack->top) {
+        data = stack->top->data;
+        if (stack->top->next) {
+            data1 = stack->top->next->data;
+        } else {
+            data1 = NULL;
+        }
+    } else {
+        data = NULL;
+        data1 = NULL;
+    }
+    goto next(data, data1, ...);
+}
+    
+__code isEmptySingleLinkedStack(struct SingleLinkedStack* stack, __code next(...), __code whenEmpty(...)) {
+    if (stack->top) {
+        goto next(...);
+    } else {
+        goto whenEmpty(...);
+    }
+}
+
+

--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/src/SpinLock.cbc	Thu Jan 16 15:12:06 2020 +0900
@@ -0,0 +1,32 @@
+#include "../context.h"
+#interface "Atomic.h"
+#interface "Lock.h"
+
+Lock* createSpinLock(struct Context* context) {
+    struct Lock* lock = new Lock();
+    struct SpinLock* spinLock = new SpinLock();
+    spinLock->lock = NULL;
+    spinLock->atomic = createAtomicReference(context);
+    lock->lock = (union Data*)spinLock;
+    lock->doLock = C_doLockSpinLock;
+    lock->doUnlock = C_doUnlockSpinLock;
+    return lock;
+}
+
+__code doLockSpinLock(struct SpinLock* lock, __code next(...)) {
+    struct Atomic* atomic = lock->atomic;
+    goto atomic->checkAndSet(&lock->lock, NULL, 1, doLockSpinLock1, doLockSpinLock);
+}
+
+__code doLockSpinLock1(struct SpinLock* lock, __code next(...)) {
+    lock->lockContext = context;
+    goto next(...);
+}
+
+__code doUnlockSpinLock(struct SpinLock* lock, __code next(...)) {
+    if (lock->lockContext == context) {
+        struct Atomic* atomic = lock->atomic;
+        goto atomic->checkAndSet(&lock->lock, 1, NULL, next(...), doUnlockSpinLock);
+    }
+    goto next(...);
+}

--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/src/Stack.h	Thu Jan 16 15:12:06 2020 +0900
@@ -0,0 +1,17 @@
+typedef struct Stack<Type, Impl>{
+        union Data* stack;
+        union Data* data;
+        union Data* data1;
+        /* Type* stack; */
+        /* Type* data; */
+        /* Type* data1; */
+        __code whenEmpty(...);
+        __code clear(Impl* stack,__code next(...));
+        __code push(Impl* stack,Type* data, __code next(...));
+        __code pop(Impl* stack, __code next(Type* data, ...));
+        __code pop2(Impl* stack, __code next(Type* data, Type* data1, ...));
+        __code isEmpty(Impl* stack, __code next(...), __code whenEmpty(...));
+        __code get(Impl* stack, __code next(Type* data, ...));
+        __code get2(Impl* stack, __code next(Type* data, Type* data1, ...));
+        __code next(...);
+} Stack;

--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/src/SynchronizedQueue.cbc	Thu Jan 16 15:12:06 2020 +0900
@@ -0,0 +1,96 @@
+#include "../context.h"
+#interface "Queue.h"
+#interface "Atomic.h"
+
+#include <stdio.h>
+
+/*
+ * Non-blocking queue of Paper: Simple, Fast, and Practical Non-Blocking and Blocking Concurrent Queue Algorithms(https://www.research.ibm.com/people/m/michael/podc-1996.pdf).
+ */
+
+Queue* createSynchronizedQueue(struct Context* context) {
+    struct Queue* queue = new Queue();
+    struct SynchronizedQueue* synchronizedQueue = new SynchronizedQueue();
+    synchronizedQueue->top = new Element(); // allocate a free node
+    synchronizedQueue->top->next = NULL;
+    synchronizedQueue->last = synchronizedQueue->top;
+    synchronizedQueue->atomic = createAtomicReference(context);
+    queue->queue = (union Data*)synchronizedQueue;
+    queue->take  = C_takeSynchronizedQueue;
+    queue->put  = C_putSynchronizedQueue;
+    queue->isEmpty = C_isEmptySynchronizedQueue;
+    queue->clear = C_clearSynchronizedQueue;
+    return queue;
+}
+
+__code clearSynchronizedQueue(struct SynchronizedQueue* queue, __code next(...)) {
+    struct Element* top = queue->top;
+    struct Atomic* atomic = queue->atomic;
+    goto atomic->checkAndSet(&queue->top, top, NULL, next(...), clearSynchronizedQueue);
+}
+
+__code putSynchronizedQueue(struct SynchronizedQueue* queue, union Data* data, __code next(...)) {
+    Element* element = new Element();
+    element->data = data;
+    element->next = NULL;
+    Element* last = queue->last;
+    Element* nextElement = last->next;
+    if (last != queue->last) {
+        goto putSynchronizedQueue();
+    }
+    if (nextElement == NULL) {
+        struct Atomic* atomic = queue->atomic;
+        goto atomic->checkAndSet(&last->next, nextElement, element, next(...), putSynchronizedQueue);
+    } else {
+        struct Atomic* atomic = queue->atomic;
+        goto atomic->checkAndSet(&queue->last, last, nextElement, putSynchronizedQueue, putSynchronizedQueue);
+    }
+}
+
+__code takeSynchronizedQueue(struct SynchronizedQueue* queue, __code next(union Data* data, ...)) {
+    struct Element* top = queue->top;
+    struct Element* last = queue->last;
+    struct Element* nextElement = top->next;
+    if (top != queue->top) {
+        goto takeSynchronizedQueue();
+    }
+    if (top == last) {
+        if (nextElement != NULL) {
+            struct Atomic* atomic = queue->atomic;
+            goto atomic->checkAndSet(&queue->last, last, nextElement, takeSynchronizedQueue, takeSynchronizedQueue);
+        }
+    } else {
+        struct Atomic* atomic = queue->atomic;
+        goto atomic->checkAndSet(&queue->top, top, nextElement, takeSynchronizedQueue1, takeSynchronizedQueue);
+    }
+    goto takeSynchronizedQueue();
+}
+
+__code takeSynchronizedQueue1(struct SynchronizedQueue* queue, __code next(union Data* data, ...), struct Element* nextElement) {
+    data = nextElement->data;
+    goto next(data, ...);
+}
+
+__code takeSynchronizedQueue1_stub(struct Context* context) {
+	SynchronizedQueue* queue = (SynchronizedQueue*)GearImpl(context, Queue, queue);
+	enum Code next = Gearef(context, Queue)->next;
+	Data** O_data = &Gearef(context, Queue)->data;
+	goto takeSynchronizedQueue1(context,
+                                queue,
+                                next,
+                                O_data,
+                                (struct Element*)Gearef(context, Atomic)->newData);
+}
+
+__code isEmptySynchronizedQueue(struct SynchronizedQueue* queue, __code next(...), __code whenEmpty(...)) {
+    struct Element* top = queue->top;
+    struct Element* last = queue->last;
+    struct Element* nextElement = top->next;
+    if (top != queue->top) {
+        goto isEmptySynchronizedQueue();
+    }
+    if (top == last && nextElement == NULL) {
+        goto whenEmpty(...);
+    }
+    goto next(...);
+}

--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/src/TaskIterator.cbc	Thu Jan 16 15:12:06 2020 +0900
@@ -0,0 +1,29 @@
+#include "../context.h"
+#interface "Iterator.h"
+#include <stdio.h>
+
+createTaskIterator(struct Context, struct List list) {
+    struct Iterator* iterator = new Iterator();
+    struct TaskIterator* taskIterator = new TaskIterator();
+    iterator->itearot = (union Data*)taskIterator;
+    iterator->exec = C_execTaskIterator;
+    iterator->barrier = C_barrierTaskIterator;
+    taskIterator->taskList = list;
+}
+
+__code execTaskIterator(struct taskIterator* iterator, struct TaskManager* taskManager, struct Context* task, __code next(...)) {
+    if (iterator->list->next == null) {
+        goto next(...);
+    }
+    iterator->list = list->next;
+    struct Context* task = (struct Context*)iterator->list->data;
+    struct TaskManager taskManager = task->taskManager;
+    taskManager->spawn(task, C_execTaskIterator);
+}
+
+__code barrierTaskIterator(struct MultiDimIterator* iterator, struct Context* task, __code next(...), __code whenWait(...)) {
+    if (__sync_fetch_and_sub(&iterator->count, 1) == 1) {
+        goto next(...);
+    }
+    goto whenWait(...);
+}

--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/src/TaskManager.h	Thu Jan 16 15:12:06 2020 +0900
@@ -0,0 +1,13 @@
+typedef struct TaskManager<Impl>{
+    union Data* taskManager;
+    struct Context* task;
+    struct Element* taskList;
+    __code spawn(Impl* taskManager, struct Context* task, __code next(...));
+    __code spawnTasks(Impl* taskManagerImpl, struct Element* taskList, __code next1(...));
+    __code setWaitTask(Impl* taskManagerImpl, struct Context* task, __code next(...));
+    __code shutdown(Impl* taskManagerImpl, __code next(...));
+    __code incrementTaskCount(Impl* taskManagerImpl, __code next(...));
+    __code decrementTaskCount(Impl* taskManagerImpl, __code next(...));
+    __code next(...);
+    __code next1(...);
+} TaskManager;

--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/src/TaskManagerImpl.cbc	Thu Jan 16 15:12:06 2020 +0900
@@ -0,0 +1,206 @@
+#include "../context.h"
+#interface "TaskManager.h"
+#interface "Iterator.h"
+#interface "Queue.h"
+#interface "Worker.h"
+
+#include <stdio.h>
+#include <unistd.h>
+
+void createWorkers(struct Context* context, TaskManagerImpl* taskManager);
+
+TaskManager* createTaskManagerImpl(struct Context* context, int numCPU, int numGPU, int numIO) {
+    struct TaskManager* taskManager = new TaskManager();
+    taskManager->spawnTasks = C_spawnTasksTaskManagerImpl;
+    taskManager->spawn = C_spawnTaskManagerImpl;
+    taskManager->shutdown  = C_shutdownTaskManagerImpl;
+    taskManager->incrementTaskCount = C_incrementTaskCountTaskManagerImpl;
+    taskManager->decrementTaskCount = C_decrementTaskCountTaskManagerImpl;
+    taskManager->setWaitTask = C_setWaitTaskTaskManagerImpl;
+    struct TaskManagerImpl* taskManagerImpl = new TaskManagerImpl();
+    // 0...numIO-1 IOProcessor
+    // numIO...numIO+numGPU-1 GPUProcessor
+    // numIO+numGPU...numIO+numGPU+numCPU-1 CPUProcessor
+    taskManagerImpl->io = 0;
+    taskManagerImpl->gpu = numIO;
+    taskManagerImpl->cpu = numIO+numGPU;
+    taskManagerImpl->maxCPU = numIO+numGPU+numCPU;
+    taskManagerImpl->numWorker = taskManagerImpl->maxCPU;
+    taskManagerImpl->sendGPUWorkerIndex = taskManagerImpl->gpu;
+    taskManagerImpl->sendCPUWorkerIndex = taskManagerImpl->cpu;
+    taskManagerImpl->taskCount = 0;
+    taskManagerImpl->loopCounter = 0;
+    createWorkers(context, taskManagerImpl);
+    taskManager->taskManager = (union Data*)taskManagerImpl;
+    return taskManager;
+}
+
+void createWorkers(struct Context* context, TaskManagerImpl* taskManager) {
+    int i = 0;
+    taskManager->workers = (Worker**)ALLOCATE_PTR_ARRAY(context, Worker, taskManager->maxCPU);
+    for (;i<taskManager->gpu;i++) {
+        Queue* queue = createSynchronizedQueue(context);
+        taskManager->workers[i] = (Worker*)createCPUWorker(context, i, queue);
+    }
+    for (;i<taskManager->cpu;i++) {
+        Queue* queue = createSynchronizedQueue(context);
+#ifdef USE_CUDAWorker
+        taskManager->workers[i] = (Worker*)createCUDAWorker(context, i, queue,0);
+#else
+        taskManager->workers[i] = (Worker*)createCPUWorker(context, i, queue);
+#endif
+    }
+    for (;i<taskManager->maxCPU;i++) {
+        Queue* queue = createSynchronizedQueue(context);
+        taskManager->workers[i] = (Worker*)createCPUWorker(context, i, queue);
+    }
+}
+
+__code spawnTasksTaskManagerImpl(struct TaskManagerImpl* taskManager, struct Element* taskList, __code next1(...)) {
+    taskManager->taskList = taskList;
+    goto spawnTasksTaskManagerImpl1();
+}
+
+__code spawnTasksTaskManagerImpl1(struct TaskManagerImpl* taskManagerImpl, struct TaskManager* taskManager) {
+    if (taskManagerImpl->taskList == NULL) {
+        goto spawnTasksTaskManagerImpl2();
+    }
+    struct Context* task = (struct Context*)taskManagerImpl->taskList->data;
+    taskManagerImpl->taskList = taskManagerImpl->taskList->next;
+    goto taskManager->setWaitTask(task, spawnTasksTaskManagerImpl1);
+}
+
+__code spawnTasksTaskManagerImpl1_stub(struct Context* context) {
+    TaskManagerImpl* taskManagerImpl = (TaskManagerImpl*)GearImpl(context, TaskManager, taskManager);
+    TaskManager* taskManager = &Gearef(context, TaskManager)->taskManager->TaskManager;
+    goto spawnTasksTaskManagerImpl1(context, taskManagerImpl, taskManager);
+}
+
+__code spawnTasksTaskManagerImpl2(struct TaskManagerImpl* taskManager, struct Element* taskList, __code next1(...)) {
+    taskManager->taskList = taskList;
+    goto spawnTasksTaskManagerImpl3();
+}
+
+__code spawnTasksTaskManagerImpl3(struct TaskManagerImpl* taskManagerImpl, __code next1(...), struct TaskManager* taskManager) {
+    if (taskManagerImpl->taskList == NULL) {
+        goto next1(...);
+    }
+    struct Context* task = (struct Context*)taskManagerImpl->taskList->data;
+    taskManagerImpl->taskList = taskManagerImpl->taskList->next;
+    goto taskManager->spawn(task, spawnTasksTaskManagerImpl3);
+}
+
+__code spawnTasksTaskManagerImpl3_stub(struct Context* context) {
+    TaskManagerImpl* taskManagerImpl = (TaskManagerImpl*)GearImpl(context, TaskManager, taskManager);
+    enum Code next1 = Gearef(context, TaskManager)->next1;
+    TaskManager* taskManager = &Gearef(context, TaskManager)->taskManager->TaskManager;
+    goto spawnTasksTaskManagerImpl3(context, taskManagerImpl, next1, taskManager);
+}
+
+__code setWaitTaskTaskManagerImpl(struct TaskManagerImpl* taskManager, struct Context* task, __code next(...)) {
+    int i = taskManager->loopCounter;
+    if (task->idg+i < task->maxIdg) {
+        struct Queue* queue = GET_WAIT_LIST(task->data[task->idg + i]);
+        taskManager->loopCounter++;
+        goto queue->put(task, setWaitTaskTaskManagerImpl);
+    }
+    taskManager->loopCounter = 0;
+    goto incrementTaskCountTaskManagerImpl();
+}
+
+__code incrementTaskCountTaskManagerImpl(struct TaskManagerImpl* taskManager, __code next(...)) {
+    __sync_fetch_and_add(&taskManager->taskCount, 1);
+    goto next(...);
+}
+
+__code decrementTaskCountTaskManagerImpl(struct TaskManagerImpl* taskManager, __code next(...)) {
+    __sync_fetch_and_sub(&taskManager->taskCount, 1);
+    goto next(...);
+}
+
+__code spawnTaskManagerImpl(struct TaskManagerImpl* taskManagerImpl, struct Context* task, __code next(...), struct TaskManager* taskManager) {
+    task->taskManager = taskManager;
+    if (task->idgCount == 0) {
+        // iterator task is normal task until spawned
+        if (task->iterator != NULL && task->iterate == 0) {
+            pthread_mutex_unlock(&taskManagerImpl->mutex);
+            struct Iterator* iterator = task->iterator;
+            goto iterator->exec(task, taskManagerImpl->cpu - taskManagerImpl->gpu, next(...));
+        }
+        goto taskSend();
+    }
+    pthread_mutex_unlock(&taskManagerImpl->mutex);
+    goto next(...);
+}
+
+__code spawnTaskManagerImpl_stub(struct Context* context) {
+    TaskManagerImpl* taskManagerImpl = (TaskManagerImpl*)GearImpl(context, TaskManager, taskManager);
+    struct Context* task = Gearef(context, TaskManager)->task;
+    TaskManager* taskManager = &Gearef(context, TaskManager)->taskManager->TaskManager;
+    goto spawnTaskManagerImpl(context,
+                              taskManagerImpl,
+                              task,
+                              Gearef(context, TaskManager)->next,
+                              taskManager);
+}
+
+
+__code taskSend(struct TaskManagerImpl* taskManager, struct Context* task, __code next(...)) {
+    // set workerId
+    if (task->gpu) {
+        goto taskSend1();
+    } else {
+        goto taskSend2();
+    }
+}
+
+__code taskSend1(struct TaskManagerImpl* taskManager, struct Context* task, __code next(...)) {
+    int workerId = taskManager->sendGPUWorkerIndex;
+    if (++taskManager->sendGPUWorkerIndex >= taskManager->cpu) {
+        taskManager->sendGPUWorkerIndex = taskManager->gpu;
+    }
+    pthread_mutex_unlock(&taskManager->mutex);
+    struct Queue* queue = taskManager->workers[workerId]->tasks;
+    goto queue->put(task, next(...));
+}
+
+__code taskSend2(struct TaskManagerImpl* taskManager, struct Context* task, __code next(...)) {
+    int workerId = taskManager->sendCPUWorkerIndex;
+    if (++taskManager->sendCPUWorkerIndex >= taskManager->maxCPU) {
+        taskManager->sendCPUWorkerIndex = taskManager->cpu;
+    }
+    pthread_mutex_unlock(&taskManager->mutex);
+    struct Queue* queue = taskManager->workers[workerId]->tasks;
+    goto queue->put(task, next(...));
+}
+
+__code shutdownTaskManagerImpl(struct TaskManagerImpl* taskManager, __code next(...)) {
+    if (taskManager->taskCount != 0) {
+        usleep(1000);
+        goto shutdownTaskManagerImpl();
+    }
+    int i = taskManager->loopCounter;
+    if (i < taskManager->numWorker) {
+        struct Queue* tasks = taskManager->workers[i]->tasks;
+        goto tasks->put(NULL, shutdownTaskManagerImpl1);
+    }
+
+    taskManager->loopCounter = 0;
+    goto shutdownTaskManagerImpl2();
+}
+
+__code shutdownTaskManagerImpl1(struct TaskManagerImpl* taskManager, __code next(...)) {
+    taskManager->loopCounter++;
+    goto shutdownTaskManagerImpl();
+}
+
+__code shutdownTaskManagerImpl2(struct TaskManagerImpl* taskManager, __code next(...)) {
+    int i = taskManager->loopCounter;
+    if (i < taskManager->numWorker) {
+        pthread_join(taskManager->workers[i]->thread, NULL);
+        taskManager->loopCounter++;
+        goto shutdownTaskManagerImpl2();
+    }
+    taskManager->loopCounter = 0;
+    goto next(...);
+}

--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/src/Timer.h	Thu Jan 16 15:12:06 2020 +0900
@@ -0,0 +1,7 @@
+typedef struct Timer<Impl>{
+        union Data* timer;
+        __code start(Impl* timer, __code next(...));
+        __code end(Impl* timer, __code next(...));
+        __code next(...);
+} Timer;
+

--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/src/TimerImpl.cbc	Thu Jan 16 15:12:06 2020 +0900
@@ -0,0 +1,30 @@
+#include <stdio.h>
+#include <sys/time.h>
+
+#include "../context.h"
+#interface "Timer.h"
+
+Timer* createTimerImpl(struct Context* context) {
+    struct Timer* timer = new Timer();
+    struct TimerImpl* timerImpl = new TimerImpl();
+    timer->timer = (union Data*)timerImpl;
+    timer->start = C_startTimer;
+    timer->end = C_endTimer;
+    return timer;
+}
+
+__code startTimer(struct TimerImpl* timer, __code next(...)) {
+    struct timeval tv;
+    gettimeofday(&tv, NULL);
+
+    timer->time = tv.tv_sec + (double)tv.tv_usec*1e-6;
+
+    goto next(...);
+}
+
+__code endTimer(struct TimerImpl* timer, __code next(...)) {
+    struct timeval tv;
+    gettimeofday(&tv, NULL);
+    printf("%0.6f\n", (tv.tv_sec+(double)tv.tv_usec*1e-6) - timer->time);
+    goto next(...);
+}

--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/src/Todo	Thu Jan 16 15:12:06 2020 +0900
@@ -0,0 +1,144 @@
+Fri May  4 20:06:48 JST 2018
+
+    par goto がある code segment は $inParGoto ではなく $hasParGoto にする
+    par goto があるばあいは goto meta ではなく goto parGotoMeta にする
+    taskList の処理も parGotoMeta で行う
+    
+    par goto の遅い理由を調べる
+        多分 Context と Synchronized Queue の生成に時間がかかってる？
+    
+    Perl スクリプトを一つにする
+    Context を生成するモジュールとstubを生成するモジュールをそれぞれオブジェクトとして作る
+    
+    Code Gear のプロトタイプを格納するオブジェクトをつくる
+
+Tue Aug  1 19:32:55 JST 2017
+ 
+    DataGear の待ち合わせ
+    DataGear の Commit
+     
+    これらは、stubとgoto meta の部分で行う
+    
+    どれに対して行うかを実行時あるいはコンパイル時に指定する必要がある
+
+    一つの解決策は、 typedefのときにannotution してあげる
+    もう一つの解決策は, Data Gear の allocation 時に指定する
+    Code Gearのプロトタイプのなかで指定する事も考えられる
+    
+    par goto時に渡す continuation で同期をとっても良い, このときにはこのcontinuation を作成するinterfaceを作る必要がある
+
+    実行時に指定してしまうと、毎回フラグのチェックが必要になる。
+    これを abstract model checking を事前に行うことで, static なコードに置き換える事はできる
+
+    例題としては, chat、dining philosophers, map reduce
+
+Fri Apr 14 18:44:09 JST 2017
+
+    struct B {
+        A a;
+        .....
+    }
+    struct A {
+        __code init(..., __code next(A a, ...));
+    }
+    par goto A->init(a);
+    // meta level
+    task->code = C_init_A;
+    task->data[idg] = ...;
+    task->data[idg + 1] = ...;
+    task->data[odg] = ...;
+    task->next = C_writeToa;
+    goto meta(context, context->TaskManager->spawn)
+
+    // lambda version?
+    par goto A->init(\A -> a = A)
+
+    // meta level
+    par goto A->init(next = \A -> a = A)
+
+Wed Mar  1 18:25:36 JST 2017
+
+    parallel_executtion/test/ を .cbc に書き直す
+    rb_tree の stub をできるだけ取り外す
+    synchornizedQueue の meta部分を分離する
+    synchronizedQueue のバグをとる
+    GPU のバグとり
+    cbc++...?
+
+Sat Jan 28 16:10:28 JST 2017
+
+    stackからpopした後、呼び出される continuation は出力を受けとる。
+    出力を受けとる stub を生成する必要がある。
+    なので、CodeGear が、そのような interface で定義されたものかどうかを調べる必要がある。
+    Stackのnext(やisEmpty)に代入された時点でわかる。なので、あまり自明な見つける方法がない。 
+    引数の異なるnextは異なる名前を持つべきか? 持たなくてもできるが...
+
+         goto next(data, ...);                                       引数で渡された continuation に移動
+         goto nodeStack->push(newNode, replaceNode1);                Interface の呼び出し。(ここで replaceNode1 が stack の戻り値を受けることがわかる。
+         goto replaceNode(traverse, traverse->current, newNode);     普通のgoto
+         goto rotateTree->next(...);                                 DataGearに格納された continuation
+
+    などをチェックする必要がある。これらの型チェックは CbC level では行われない。(CbCはmeta levelだから)
+
+     戻り値の部分は interface に記述させるという手もあるな。
+
+
+Sun Jan 22 20:11:28 JST 2017
+
+    TaskManagerから必要なCPUWorkerを生成する
+    WorkerはcreateWorker時に新しくthreadを作る
+    
+    TaskManager->createTaskで新しいContextを生成する
+    この時点でWorkerを番号で指定する
+    このContextにGearefで値を設定していく
+    待ち合わせ用のDSを設定する
+    taskManager->spawnでWorkerにcontextを送る
+
+Fri Jan 13 17:47:40 JST 2017
+
+    Task は contextを直接使うことにする
+        DS には, まっているcontextをListを作る
+        context に実行中断中のCS の番号をいれるフィールドを用意する
+        待っているDS のcount
+    createTaskの手順
+        新しくcontextを作る
+            allocate 用のheap も用意
+            もとのcontextを全部copyする or 必要なものだけcopyする
+            待ち合わせのDS群を指定する
+            終わったあとの行き先を指定する(default は task_exit)
+            exception の行き先も必要な指定する
+            待っているDSが全部揃っていたら active Queueに入れる
+    task の実行
+        taskの実行後、 goto meta する直前で code gear commit を呼んで, Reader list を消化する
+        複数から参照されるDSは一旦localに書き出して, その後atomic に書き出す
+        複数から参照されるDSは何かしら宣言が必要
+            つまり DS には 一つ一つ owner がいる
+
+Mon Nov 28 17:39:39 JST 2016
+
+    Task,TaskManager,Workerのインターフェースの実装を作成する
+    Taskを一旦Treeに入れずに直接Queueに入れる
+
+    Task
+        CodeGen
+            IDataSeg
+            IDataSeg
+            ...
+        idsCount
+        nextTask(can be C_exit)
+            ODataSeg?
+
+    TaskManager
+        createWorker
+        spawn (any,cpu,GPU)
+        taskSend
+        activeQueue
+        shutdown
+        deadlockDetectid
+
+    SynchronizedQueue * Workerの数だけ
+
+    Worker
+        execute
+        taskRecive
+        shutdown

--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/src/Tree.h	Thu Jan 16 15:12:06 2020 +0900
@@ -0,0 +1,12 @@
+typedef struct Tree<Type, Impl>{
+    /* future Code */
+    /* Type* tree; */
+    /* Type* node; */
+    union Data* tree;
+    struct Node* node;
+    __code put(Impl* tree,Type* node, __code next(...));
+    // __code get(Impl* tree, __code next(...));
+    __code remove(Impl* tree,Type* node, __code next(...));
+    // __code clearRedBlackTree();
+    __code next(...);
+} Tree;

--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/src/Worker.h	Thu Jan 16 15:12:06 2020 +0900
@@ -0,0 +1,10 @@
+typedef struct Worker<Impl>{
+    union Data* worker;
+    struct Queue* tasks;
+    struct Context* task;
+    pthread_t thread;
+    struct TaskManager* taskManager;
+    __code taskReceive(Impl* worker, struct Queue* tasks);
+    __code shutdown(Impl* worker);
+    __code next(...);
+} Worker;

--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/src/auto_generate_context.pl	Thu Jan 16 15:12:06 2020 +0900
@@ -0,0 +1,87 @@
+#!/usr/bin/env perl
+use strict;
+use warnings;
+use FindBin;
+use lib "$FindBin::Bin/lib";
+
+use Gears::Context;
+use Getopt::Std;
+my %opt;
+getopts("w" => \%opt);
+
+my $h = find_target_from_camke_list();
+
+my $target = shift;
+unless (exists $h->{$target}) {
+  map { print "$_\n" }  keys %$h;
+  print "invalid target name\n";
+  exit 0;
+}
+
+my @cbc_files;
+
+for my $file_name (split / /, $h->{$target}) {
+  if ($file_name =~ /^\s*$/) {
+    next;
+  }
+  chomp $file_name;
+  push(@cbc_files,$file_name);
+}
+
+my $output     = $opt{w} ? "$FindBin::Bin/context.h" : "stdout";
+my $gears      = Gears::Context->new(compile_sources => \@cbc_files, find_root => $FindBin::Bin, output => $output);
+my $data_gears = $gears->extraction_dg_compile_sources();
+my $g          = $gears->set_data_gear_header_path();
+
+my $dg2path    = $gears->update_dg_each_header_path($data_gears,$g);
+
+my $tree = $gears->createImplTree_from_header($dg2path);
+$gears->tree2create_context_h($tree);
+
+
+sub find_target_from_camke_list {
+  open my $fh, '<', "CMakeLists.txt";
+  my $in_gears = 0;
+  my $target;
+  my %res;
+
+  while (my $line = <$fh>) {
+    if ($in_gears == 1) {
+      $in_gears++;
+      next;
+    }
+
+    if ($in_gears == 2) {
+       $line =~ s/\s*(\w+)\s*/$1/g;
+       $target = $line;
+       chomp $target;
+       $in_gears++;
+       next;
+    }
+
+    if ($in_gears == 3) {
+      $in_gears++;
+      next;
+    }
+
+    if ($in_gears == 4) {
+      $res{$target} = $line;
+      chomp $res{$target};
+      $in_gears = 0;
+      next;
+    }
+
+    if ($line =~ /^GearsCommand\(/) {
+        $in_gears++;
+    }
+
+  }
+  return \%res;
+}
+
+#GearsCommand(
+#  TARGET
+#      rbtree
+#  SOURCES
+#      SingleLinkedQueue.cbc test/rbTree_test.cbc RedBlackTree.cbc SingleLinkedStack.cbc compare.c
+#)

--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/src/compare.c	Thu Jan 16 15:12:06 2020 +0900
@@ -0,0 +1,13 @@
+#include "context.h"
+
+enum Relational compare(struct Node* node1, struct Node* node2) {
+    int key1 = node1->key;
+    int key2 = node2->key;
+    if (key1 == key2) {
+        return EQ;
+    } else if (key1 < key2) {
+        return GT;
+    } else {
+        return LT;
+    }
+}

--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/src/cuda.c	Thu Jan 16 15:12:06 2020 +0900
@@ -0,0 +1,95 @@
+#include <stdio.h>
+#include <sys/time.h>
+#include <string.h>
+#include <stdlib.h>
+
+// includes, project
+#include <driver_types.h>
+#include <cuda_runtime.h>
+#include <cuda.h>
+#include "helper_cuda.h"
+#include "pthread.h"
+
+#include "context.h"
+
+/*
+struct Context {
+    int next;
+    struct Worker* worker;
+    struct TaskManager* taskManager;
+    int codeNum;
+    void  (**code) (struct Context*);
+    void* heapStart;
+    void* heap;
+    long heapLimit;
+    int dataNum;
+    int idgCount; //number of waiting dataGear
+    int idg;
+    int maxIdg;
+    int odg;
+    int maxOdg;
+    int workerId;
+    struct Context* task;
+    struct Queue* tasks;
+    int num_exec;
+    CUmodule module;
+    CUfunction function;
+    union Data **data;
+
+    // multi dimension parameter
+    int iterate;
+    struct Iterator* iterator;
+};
+
+struct CUDAWorker {
+    CUdevice device;
+    CUcontext cuCtx;
+    pthread_t thread;
+    struct Context* context;
+    int id;
+    struct Queue* tasks;
+    int runFlag;
+    int next;
+    int numStream;
+    CUstream *stream;
+} CUDAWorker;
+
+struct LoopCounter {
+    int i;
+} LoopCounter;
+
+struct Array {
+    int size;
+    int index;
+    int prefix;
+    int* array;
+} Array;
+*/
+
+void cudaInit(struct CUDAWorker *cudaWorker,int phase, int deviceNum) {
+    // initialize and load kernel
+    cudaWorker->numStream = 1; // number of stream
+    //    cudaWorker->stream = NEWN(cudaWorker->numStream, CUstream );
+    if (phase==0)
+        checkCudaErrors(cuInit(0));
+    if (phase==0)
+        checkCudaErrors(cuDeviceGet(&cudaWorker->device, deviceNum));
+    if (phase==0)
+        checkCudaErrors(cuCtxCreate(&cudaWorker->cuCtx, CU_CTX_SCHED_SPIN, cudaWorker->device));
+    //    if (cudaWorker->num_stream) {
+    //        for (int i=0;i<cudaWorker->num_stream;i++)
+    //            checkCudaErrors(cuStreamCreate(&cudaWorker->stream[i],0));
+    //    }
+    printf("cuda Init: Done\n");
+}
+
+void cudaLoadFunction(struct Context* context, char* filename, char* function) {
+    checkCudaErrors(cuModuleLoad(&context->module, filename));
+    checkCudaErrors(cuModuleGetFunction(&context->function, context->module, function));
+}
+
+void cudaShutdown(struct CUDAWorker *worker) {
+    //    for (int i=0;i<worker->num_stream;i++)
+    //        checkCudaErrors(cuStreamDestroy(worker->stream[i]));
+    checkCudaErrors(cuCtxDestroy(worker->cuCtx));
+}

--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/src/examples/bitonicSort/CUDAbitonicSwap.cu	Thu Jan 16 15:12:06 2020 +0900
@@ -0,0 +1,32 @@
+extern "C" {
+    struct Integer {
+        int value;
+    };
+    struct SortArray {
+        struct Integer *array;
+        int loopCounter;
+        int block;
+        int first;
+        int prefix;
+    };
+    __global__ void bitonicSwap(struct Integer* array, struct SortArray* sortArray) {
+        int block = sortArray->block;
+        int first = sortArray->first;
+        int prefix = sortArray->prefix;
+        int i = 0;
+C_bitonicSwap:
+        if (i < prefix) {
+            int index = i + (blockIdx.x * blockDim.x + threadIdx.x) * prefix;
+            int position = index/block;
+            int index1 = index+block*position;
+            int index2 = (first == 1)? ((block<<1)*(position+1))-(index1%block)-1 : index1+block;
+            if (array[index2].value < array[index1].value) {
+                struct Integer tmp = array[index1];
+                array[index1] = array[index2];
+                array[index2] = tmp;
+            }
+            i++;
+            goto C_bitonicSwap;
+        }
+    }
+}

--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/src/examples/bitonicSort/SortArray.h	Thu Jan 16 15:12:06 2020 +0900
@@ -0,0 +1,8 @@
+typedef struct SortArray<Impl>{
+    union Data* sortArray;
+    struct Integer *array;
+    int loopCounter;
+    int block;
+    int first;
+    int prefix;
+} SortArray;

--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/src/examples/bitonicSort/bitonicSort.cbc	Thu Jan 16 15:12:06 2020 +0900
@@ -0,0 +1,89 @@
+#include <stdio.h>
+#include <string.h>
+#include <stdlib.h>
+#include <unistd.h>
+#include <math.h>
+
+#include "../../../context.h"
+#interface "TaskManager.h"
+
+int cpu_num = 1;
+int length = 1024;
+int split  = 8;
+int gpu_num = 0;
+int CPU_ANY = -1;
+int CPU_CUDA = -1;
+
+__code initDataGears(struct LoopCounter* loopCounter, struct TaskManager* taskManager) {
+    // loopCounter->tree = createRedBlackTree(context);
+    loopCounter->i = 0;
+    taskManager->taskManager = (union Data*)createTaskManagerImpl(context, cpu_num, gpu_num, 0);
+    goto code1();
+}
+
+__code code1(struct LoopCounter* loopCounter) {
+    printf("cpus:\t\t%d\n", cpu_num);
+    printf("gpus:\t\t%d\n", gpu_num);
+    printf("length:\t\t%d\n", length);
+    printf("length/task:\t%d\n", length/split);
+    /* puts("queue"); */
+    /* print_queue(context->data[ActiveQueue]->queue.first); */
+    /* puts("tree"); */
+    /* print_tree(context->data[Tree]->tree.root); */
+    /* puts("result"); */
+
+    goto createTask1();
+}
+
+__code createTask1(struct LoopCounter* loopCounter, struct TaskManager* taskManager) {
+    struct SortArray* outputSortArray = new SortArray();
+    struct SortArray* inputSortArray = outputSortArray;
+    struct Timer* timer = createTimerImpl(context);
+
+    par goto makeArray(outputSortArray, timer, __exit);
+
+    for (int i=2; i <= length; i=2*i) {
+        int first = 1;
+        for (int j=i>>1; j > 0; j=j>>1) {
+            outputSortArray = new SortArray();
+            inputSortArray->prefix = length/2/split;
+            inputSortArray->block = j;
+            inputSortArray->first = first;
+            par goto bitonicSwap(inputSortArray, outputSortArray, iterate(split), __exit);
+            first = 0;
+            inputSortArray = outputSortArray;
+        }
+    }
+
+    par goto printArray(inputSortArray, timer, __exit);
+
+    goto code2();
+}
+
+__code code2(struct TaskManager* taskManager) {
+    goto taskManager->shutdown(exit_code);
+}
+
+__code code2_stub(struct Context* context) {
+    goto code2(context, &Gearef(context, TaskManager)->taskManager->TaskManager);
+}
+
+void init(int argc, char** argv) {
+    for (int i = 1; argv[i]; ++i) {
+        if (strcmp(argv[i], "-cpu") == 0)
+            cpu_num = (int)atoi(argv[i+1]);
+        else if (strcmp(argv[i], "-l") == 0)
+            length = (int)atoi(argv[i+1]);
+        else if (strcmp(argv[i], "-s") == 0)
+            split = (int)atoi(argv[i+1]);
+        else if (strcmp(argv[i], "-cuda") == 0) {
+            gpu_num = 1;
+            CPU_CUDA = 0;
+        }
+    }
+}
+
+int main(int argc, char** argv) {
+    init(argc, argv);
+    goto initDataGears();
+}

--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/src/examples/bitonicSort/bitonicSwap.cbc	Thu Jan 16 15:12:06 2020 +0900
@@ -0,0 +1,60 @@
+#include "../../../context.h"
+#include <stdio.h>
+
+#ifdef USE_CUDAWorker
+extern void cudaLoadFunction(struct Context* context, char* filename, char* function);
+#endif
+
+__code bitonicSwap(struct SortArray* inputArray, struct MultiDim* multiDim, __code next(struct SortArray* output, ...), struct LoopCounter* loopCounter) {
+    int block = inputArray->block;
+    int first = inputArray->first;
+    if (loopCounter->i < inputArray->prefix) {
+        int index = loopCounter->i + multiDim->x * inputArray->prefix;
+        int position = index/block;
+        int index1 = index+block*position;
+        int index2 = (first == 1)? ((block<<1)*(position+1))-(index1%block)-1 : index1+block;
+        struct Integer* array = inputArray->array;
+        if (array[index2].value < array[index1].value) {
+            struct Integer tmp = array[index1];
+            array[index1] = array[index2];
+            array[index2] = tmp;
+        }
+        loopCounter->i++;
+        goto bitonicSwap();
+    }
+    loopCounter->i = 0;
+    output->array = inputArray->array;
+    goto next(output, ...);
+}
+
+__code bitonicSwap_stub(struct Context* context) {
+#ifdef USE_CUDAWorker
+    if (context->gpu) {
+        SortArray* inputSortArray  = &context->data[context->idg]->SortArray;
+        SortArray* outputSortArray = &context->data[context->odg]->SortArray;
+        CUDABuffer* buffer = &ALLOCATE(context, CUDABuffer)->CUDABuffer;
+        buffer->inputData = (union Data**)ALLOCATE_PTR_ARRAY(context, SortArray, 2);
+        buffer->inputData[0] = (union Data*)inputSortArray->array;
+        buffer->inputData[1] = (union Data*)inputSortArray;
+        buffer->outputData = NULL;
+        buffer->inputLen = 2;
+        buffer->outputLen = 0;
+        //continuationにそってGPUworkerに戻る
+        outputSortArray->array = inputSortArray->array;
+        Executor* executor = context->worker->worker->CUDAWorker.executor;
+        executor->executor->CUDAExecutor.buffer = buffer;
+        cudaLoadFunction(context, "c/examples/bitonicSort/CUDAbitonicSwap.ptx", "bitonicSwap");
+        Gearef(context, Executor)->executor = (union Data*)executor;
+        Gearef(context, Executor)->task = context;
+        Gearef(context, Executor)->next = context->next;
+        goto meta(context, executor->read);
+    }
+#endif
+    SortArray** O_output = (struct SortArray **)&context->data[context->odg];
+    goto bitonicSwap(context,
+            &context->data[context->idg]->SortArray,
+            &context->data[context->idg+1]->MultiDim,
+            context->next,
+            O_output,
+            Gearef(context, LoopCounter));
+}

--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/src/examples/bitonicSort/makeArray.cbc	Thu Jan 16 15:12:06 2020 +0900
@@ -0,0 +1,20 @@
+#include "../../../context.h"
+#include <stdio.h>
+#interface "Timer.h"
+
+extern int length;
+__code makeArray(__code next(struct SortArray* output, struct Timer* output1, ...)){
+    if (output->loopCounter == 0){
+        output->array = (Integer*)ALLOCATE_ARRAY(context, Integer, length);
+        srand((unsigned) time(NULL));
+    }
+    if (output->loopCounter == GET_LEN(output->array)){
+        printf("created Array\n");
+        output->loopCounter = 0;
+        goto output1->start(next(...));
+    }
+    output->array[output->loopCounter].value = rand() % 1000;
+    //printf("%d\n", output->array[output->loopCounter]->value);
+    output->loopCounter++;
+    goto makeArray();
+}