# HG changeset patch
# User Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
# Date 1504365676 -32400
# Node ID c5cd9888bf2aea0bb0c3b323bf89538b4b248f4b
# Parent  83c9aeb1fe3e1a72813fc66fa707101167ba3228
Fix bitonicSort

diff -r 83c9aeb1fe3e -r c5cd9888bf2a src/parallel_execution/CMakeLists.txt
--- a/src/parallel_execution/CMakeLists.txt	Fri Sep 01 06:41:02 2017 +0900
+++ b/src/parallel_execution/CMakeLists.txt	Sun Sep 03 00:21:16 2017 +0900
@@ -73,7 +73,7 @@
   TARGET
       bitonicSort
   SOURCES
-      examples/bitonicSort/bitonicSort.cbc examples/bitonicSort/swap.cbc CPUWorker.cbc TaskManagerImpl.cbc SingleLinkedQueue.cbc SynchronizedQueue.cbc MultiDimIterator.cbc TimeImpl.cbc
+      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 TimeImpl.cbc
 )
 
 if (${USE_CUDA})
diff -r 83c9aeb1fe3e -r c5cd9888bf2a src/parallel_execution/context.h
--- a/src/parallel_execution/context.h	Fri Sep 01 06:41:02 2017 +0900
+++ b/src/parallel_execution/context.h	Sun Sep 03 00:21:16 2017 +0900
@@ -34,6 +34,15 @@
     meta->size = 1;     \
     data; })
 
+#define ALLOCATE_ARRAY(context, t, len) ({ \
+    struct Meta* meta = (struct Meta*)context->heap;\
+    context->heap += sizeof(struct Meta);\
+    union Data* data = context->heap; \
+    context->heap += sizeof(struct t)*len; \
+    meta->type = D_##t; \
+    meta->size = len;     \
+    data;   })
+
 #define ALLOC(context, t) (&ALLOCATE(context, t)->t)
 
 #define ALLOC_ARRAY(context, dseg, len) ({\
@@ -294,20 +303,12 @@
     struct Integer {
         int value;
     } Integer;
-    struct SortArray {//そもそもこれは必要なのか？
-        struct SortArray *sortArray;
-        struct Integer **array;//Array arrayじゃできない？
+    struct SortArray {
+        struct Integer *array; //Array arrayじゃできない？
         int loopCounter;
-        int loopCounter2;
-        int loopCounter3;
-        int sort_finish;
+        int block;
+        int first;
         int prefix;
-        enum Code make_array;
-        enum Code print;
-        enum Code bitonic_sort;
-        enum Code kernel;
-        enum Code kernel2;
-        enum Code swap;
     } SortArray;
     struct Iterator {
         union Data* iterator;
diff -r 83c9aeb1fe3e -r c5cd9888bf2a src/parallel_execution/cuda.c
--- a/src/parallel_execution/cuda.c	Fri Sep 01 06:41:02 2017 +0900
+++ b/src/parallel_execution/cuda.c	Sun Sep 03 00:21:16 2017 +0900
@@ -101,7 +101,7 @@
      checkCudaErrors(cuLaunchKernel(context->function,
                        1, 1, 1,
                        1, 1, 1,
-                                 0, NULL , args, NULL));
+                       0, NULL, args, NULL));
 
     //結果を取ってくるコマンドを入力する
     //コマンドの終了待ちを行う   
diff -r 83c9aeb1fe3e -r c5cd9888bf2a src/parallel_execution/examples/bitonicSort/bitonicSort.cbc
--- a/src/parallel_execution/examples/bitonicSort/bitonicSort.cbc	Fri Sep 01 06:41:02 2017 +0900
+++ b/src/parallel_execution/examples/bitonicSort/bitonicSort.cbc	Sun Sep 03 00:21:16 2017 +0900
@@ -71,77 +71,28 @@
 }
 
 __code createTask1(struct LoopCounter* loopCounter, struct TaskManager* taskManager) {
-    int logN = (int)log2((float)length);
-    struct Context** tasks = (struct Context**)ALLOC_ARRAY(context, Context, logN*(1+logN)/2 + 2);
-    int taskCount = 0;
-
     struct SortArray* outputSortArray = &ALLOCATE_DATA_GEAR(context, SortArray)->SortArray;
     struct SortArray* inputSortArray = outputSortArray;
     struct Time* time = createTimeImpl(context);
 
-    // par goto makeArray(outputSortArray, time, _exit);
-    // par goto makeArray(outputSortArray, time, time->start); ?
-    struct Context* task = NEW(struct Context);
-    initContext(task);
-    task->next = C_makeArray;
-    task->idgCount = 0;
-    task->idg = task->dataNum;
-    task->maxIdg = task->idg;
-    task->odg = task->maxIdg;
-    task->data[task->odg] = (union Data*)outputSortArray;
-    task->data[task->odg+1] = (union Data*)time;
-    task->maxOdg = task->odg + 2;
-    tasks[taskCount] = task;
+    par goto makeArray(outputSortArray, time, _exit);
 
-    taskCount++;
     for (int i=2; i <= length; i=2*i) {
         int first = 1;
         for (int j=i>>1; j > 0; j=j>>1) {
             outputSortArray = &ALLOCATE_DATA_GEAR(context, SortArray)->SortArray;
-            struct Context* task = NEW(struct Context);
-            initContext(task);
-            struct Integer* integer1 = &ALLOCATE_DATA_GEAR(context, Integer)->Integer;
-            struct Integer* integer2 = &ALLOCATE_DATA_GEAR(context, Integer)->Integer;
-            integer1->value = j;
-            integer2->value = first;
             inputSortArray->prefix = length/2/split;
-            task->next = C_bitonicSwap;
-            task->iterator = createOneDimIterator(context, split);
-            task->idgCount = 1;
-            task->idg = task->dataNum;
-            task->data[task->idg] = (union Data*)inputSortArray;
-            task->data[task->idg+1] = (union Data*)integer1;
-            task->data[task->idg+2] = (union Data*)integer2;
-            task->maxIdg = task->idg + 3;
-            task->odg = task->maxIdg;
-            task->data[task->odg] = (union Data*)outputSortArray;
-            task->maxOdg = task->odg + 1;
-            tasks[taskCount] = task;
-            taskCount++;
+            inputSortArray->block = j;
+            inputSortArray->first = first;
+            par goto bitonicSwap(inputSortArray, outputSortArray, iterate(split), _exit);
             first = 0;
             inputSortArray = outputSortArray;
         }
     }
 
-    // par goto printArray(inputSortArray, time, __exit)
-    // par goto printArray(inputSortArray, time, time->exit)?
-    task = NEW(struct Context);
-    initContext(task);
-    task->next = C_printArray;
-    task->idgCount = 2;
-    task->idg = task->dataNum;
-    task->data[task->idg] = (union Data*)inputSortArray;
-    task->data[task->idg+1] = (union Data*)time;
-    task->maxIdg = task->idg + 2;
-    task->odg = task->maxIdg;
-    task->maxOdg = task->odg;
-    tasks[taskCount] = task;
+    par goto printArray(inputSortArray, time, __exit);
 
-    taskManager->contexts = tasks;
-    // goto code2();
-    taskManager->next1 = C_code2;
-    goto meta(context, taskManager->taskManager->TaskManager.spawnTasks);
-    //goto meta(context, taskManager->taskManager->TaskManager.shutdown);
+    goto code2();
 }
 
 __code code2(struct LoopCounter* loopCounter, struct TaskManager* taskManager, struct Time* time) {
@@ -173,73 +124,3 @@
 
     goto start_code(main_context);
 }
-
-__code makeArray(struct Time* timeInterface, __code next(struct SortArray* output, struct Time* output1, ...)){
-    struct SortArray* output = *O_output;
-    struct Time* output1 = *O_output1;
-    if (output->loopCounter == 0){
-        output->array = (Integer**)ALLOC_ARRAY(context, Integer, length);
-        srand((unsigned) time(NULL));
-    }
-    if (output->loopCounter == GET_SIZE(output->array)){
-        printf("created Array\n");
-        output->loopCounter = 0;
-        timeInterface->time = (union Data*)output1;
-        timeInterface->next = context->next;
-        *O_output = output;
-        *O_output1 = output1;
-        goto meta(context, output1->start);
-    }
-    struct Integer* integer = new Integer();
-    integer->value = rand() % 1000;
-    output->array[output->loopCounter] = integer;
-    //printf("%d\n", output->array[output->loopCounter]->value);
-    output->loopCounter++;
-    *O_output = output;
-    *O_output1 = output1;
-    goto meta(context, C_makeArray);
-}
-
-__code makeArray_stub(struct Context* context) {
-    SortArray** O_output = (struct SortArray**)&context->data[context->odg];
-    Time** O_output1 = (struct Time**)&context->data[context->odg+1];
-    goto makeArray(context,
-                   Gearef(context, Time),
-                   context->next,
-                   O_output,
-                   O_output1);
-}
-
-__code printArray(struct Time* timeInterface, struct SortArray* inputArray, struct Time* inputTime, __code next(...)){
-    timeInterface->time = (union Data*)inputTime;
-    timeInterface->next = C_printArray1;
-    goto meta(context, inputTime->end);
-}
-
-__code printArray_stub(struct Context* context) {
-    goto printArray(context,
-                   Gearef(context, Time),
-                   &context->data[context->idg]->SortArray,
-                   &context->data[context->idg+1]->Time,
-                   context->next);
-}
-
-__code printArray1(struct SortArray* inputArray, __code next(...)){
-    //printf("%d\n", inputArray->array[inputArray->loopCounter]->value);
-    inputArray->loopCounter++;
-    if (inputArray->loopCounter == GET_SIZE(inputArray->array)){
-        inputArray->loopCounter = 0;
-        goto meta(context, next);
-    }
-    if (inputArray->array[inputArray->loopCounter-1]->value > inputArray->array[inputArray->loopCounter]->value) {
-        printf("wrong result\n");
-        goto meta(context, next);
-    }
-    goto meta(context, C_printArray1);
-}
-
-__code printArray1_stub(struct Context* context) {
-    goto printArray1(context,
-                   &context->data[context->idg]->SortArray,
-                   context->next);
-}
diff -r 83c9aeb1fe3e -r c5cd9888bf2a src/parallel_execution/examples/bitonicSort/bitonicSwap.cbc
--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/src/parallel_execution/examples/bitonicSort/bitonicSwap.cbc	Sun Sep 03 00:21:16 2017 +0900
@@ -0,0 +1,36 @@
+#include "../../../context.h"
+#include <stdio.h>
+
+__code bitonicSwap(struct SortArray* inputArray, struct MultiDim* multiDim, __code next(struct SortArray* output, ...), struct LoopCounter* loopCounter) {
+    struct SortArray* output = *O_output;
+    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 meta(context, C_bitonicSwap);
+    }
+    loopCounter->i = 0;
+    output->array = inputArray->array;
+    *O_output = output;
+    goto meta(context, next);
+}
+
+__code bitonicSwap_stub(struct Context* context) {
+    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));
+}
diff -r 83c9aeb1fe3e -r c5cd9888bf2a src/parallel_execution/examples/bitonicSort/makeArray.cbc
--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/src/parallel_execution/examples/bitonicSort/makeArray.cbc	Sun Sep 03 00:21:16 2017 +0900
@@ -0,0 +1,35 @@
+#include "../../../context.h"
+#include <stdio.h>
+extern int length;
+__code makeArray(__code next(struct SortArray* output, struct Time* output1, ...)){
+    struct SortArray* output = *O_output;
+    struct Time* output1 = *O_output1;
+    if (output->loopCounter == 0){
+        output->array = (Integer*)ALLOCATE_ARRAY(context, Integer, length);
+        srand((unsigned) time(NULL));
+    }
+    if (output->loopCounter == GET_SIZE(output->array)){
+        printf("created Array\n");
+        output->loopCounter = 0;
+        Gearef(context, Time)->time = (union Data*)output1;
+        Gearef(context, Time)->next = context->next;
+        *O_output = output;
+        *O_output1 = output1;
+        goto meta(context, output1->start);
+    }
+    output->array[output->loopCounter].value = rand() % 1000;
+    //printf("%d\n", output->array[output->loopCounter]->value);
+    output->loopCounter++;
+    *O_output = output;
+    *O_output1 = output1;
+    goto meta(context, C_makeArray);
+}
+
+__code makeArray_stub(struct Context* context) {
+    SortArray** O_output = (struct SortArray**)&context->data[context->odg];
+    Time** O_output1 = (struct Time**)&context->data[context->odg+1];
+    goto makeArray(context,
+                   context->next,
+                   O_output,
+                   O_output1);
+}
diff -r 83c9aeb1fe3e -r c5cd9888bf2a src/parallel_execution/examples/bitonicSort/printArray.cbc
--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/src/parallel_execution/examples/bitonicSort/printArray.cbc	Sun Sep 03 00:21:16 2017 +0900
@@ -0,0 +1,36 @@
+#include "../../../context.h"
+#include <stdio.h>
+
+__code printArray(struct SortArray* inputArray, struct Time* inputTime, __code next(...)){
+    Gearef(context, Time)->time = (union Data*)inputTime;
+    Gearef(context, Time)->next = C_printArray1;
+    goto meta(context, inputTime->end);
+}
+
+__code printArray_stub(struct Context* context) {
+    goto printArray(context,
+                   &context->data[context->idg]->SortArray,
+                   &context->data[context->idg+1]->Time,
+                   context->next);
+}
+
+__code printArray1(struct SortArray* inputArray, __code next(...)){
+    //printf("%d\n", inputArray->array[inputArray->loopCounter].value);
+    inputArray->loopCounter++;
+    if (inputArray->loopCounter == GET_SIZE(inputArray->array)){
+        printf("sort completed\n");
+        inputArray->loopCounter = 0;
+        goto meta(context, next);
+    }
+    if (inputArray->array[inputArray->loopCounter-1].value > inputArray->array[inputArray->loopCounter].value) {
+        printf("wrong result\n");
+        goto meta(context, next);
+    }
+    goto meta(context, C_printArray1);
+}
+
+__code printArray1_stub(struct Context* context) {
+    goto printArray1(context,
+                   &context->data[context->idg]->SortArray,
+                   context->next);
+}
diff -r 83c9aeb1fe3e -r c5cd9888bf2a src/parallel_execution/examples/bitonicSort/swap.cbc
--- a/src/parallel_execution/examples/bitonicSort/swap.cbc	Fri Sep 01 06:41:02 2017 +0900
+++ /dev/null	Thu Jan 01 00:00:00 1970 +0000
@@ -1,36 +0,0 @@
-#include "../../../context.h"
-#include <stdio.h>
-
-__code bitonicSwap(struct SortArray* inputArray, struct Integer* block, struct Integer* first, struct Integer* i, struct LoopCounter* loopCounter, __code next(struct SortArray* output, ...)) {
-    struct SortArray* output = *O_output;
-    if (loopCounter->i < inputArray->prefix) {
-        int index = loopCounter->i + i->value * inputArray->prefix;
-        int position = index/block->value;
-        int index1 = index+block->value*position;
-        int index2 = (first->value == 1)? ((block->value<<1)*(position+1))-(index1%block->value)-1 : index1+block->value;
-        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 meta(context, C_bitonicSwap);
-    }
-    loopCounter->i = 0;
-    output->array = inputArray->array;
-    *O_output = output;
-    goto meta(context, next);
-}
-
-__code bitonicSwap_stub(struct Context* context) {
-    SortArray** O_output = (struct SortArray **)&context->data[context->odg];
-    goto bitonicSwap(context,
-                     &context->data[context->idg]->SortArray,
-                     &context->data[context->idg+1]->Integer,
-                     &context->data[context->idg+2]->Integer,
-                     &context->data[context->idg+3]->Integer,
-                     Gearef(context, LoopCounter),
-                     context->next,
-                     O_output);
-}
diff -r 83c9aeb1fe3e -r c5cd9888bf2a src/parallel_execution/generate_stub.pl
--- a/src/parallel_execution/generate_stub.pl	Fri Sep 01 06:41:02 2017 +0900
+++ b/src/parallel_execution/generate_stub.pl	Sun Sep 03 00:21:16 2017 +0900
@@ -178,15 +178,14 @@
             while($args) {
                 if ($args =~ s/(^\s*,\s*)//) {
                 }
-                if ($args =~ s/^(\s)*\_\_code\s+(\w+)\(//) {
+                if ($args =~ s/^(\s)*\_\_code\s+(\w+)\((.*?)\)//) {
                     $inputIncFlag = 0;
+                    $outputCount = split(/,/,$3);
+                    $outputCount--;
                 } elsif ($args =~ s/^(struct|union) (\w+)(\*)+\s(\w+)//) {
                     if($inputIncFlag) {
                         $inputCount++;
                     }
-                    else {
-                        $outputCount++;
-                    }
                 } elsif ($args =~ s/(.*,)//) {
                 } else {
                     last;
@@ -435,25 +434,25 @@
                     print $fd "${prev}struct Context* task;\n";
                 }
                 my $initTask = << "EOFEOF";
-${prev}task = NEW(struct Context);
-${prev}initContext(task);
-${prev}task->next = C_$codeGearName;
-${prev}task->idgCount = $inputCount;
-${prev}task->idg = task->dataNum;
-${prev}task->maxIdg = task->idg + $inputCount;
-${prev}task->odg = task->maxIdg;
-${prev}task->maxOdg = task->odg + $outputCount;
+                ${prev}task = NEW(struct Context);
+                ${prev}initContext(task);
+                ${prev}task->next = C_$codeGearName;
+                ${prev}task->idgCount = $inputCount;
+                ${prev}task->idg = task->dataNum;
+                ${prev}task->maxIdg = task->idg + $inputCount;
+                ${prev}task->odg = task->maxIdg;
+                ${prev}task->maxOdg = task->odg + $outputCount;
 EOFEOF
                 print $fd $initTask;
                 if (@iterateCounts) {
                     print $fd "${prev}task->iterate = 0;\n";
                     my $len = @iterateCounts;
                     if ($len == 1) {
-                      print $fd "${prev}task->iterator = createMultiDimIterator(context, $iterateCounts[0], 1, 1);\n";
+                        print $fd "${prev}task->iterator = createMultiDimIterator(context, $iterateCounts[0], 1, 1);\n";
                     } elsif ($len == 2) {
-                      print $fd "${prev}task->iterator = createMultiDimIterator(context, $iterateCounts[0], $iterateCounts[1], 1);\n";
+                        print $fd "${prev}task->iterator = createMultiDimIterator(context, $iterateCounts[0], $iterateCounts[1], 1);\n";
                     } elsif ($len == 3) {
-                      print $fd "${prev}task->iterator = createMultiDimIterator(context, $iterateCounts[0], $iterateCounts[1], $iterateCounts[2]);\n";
+                        print $fd "${prev}task->iterator = createMultiDimIterator(context, $iterateCounts[0], $iterateCounts[1], $iterateCounts[2]);\n";
                     }
                 }
                 for my $i (0..$inputCount-1) {
@@ -464,11 +463,11 @@
                     print $fd "${prev}task->data[task->odg+$i] = (union Data*)@dataGears[$inputCount+$i];\n";
                 }
                 my $putTask = << "EOFEOF";
-${prev}element = &ALLOCATE(context, Element)->Element;
-${prev}element->next = NULL;
-${prev}element->data = (union Data*)task;
-${prev}queue->last->next  = element;
-${prev}queue->last = element;
+                ${prev}element = &ALLOCATE(context, Element)->Element;
+                ${prev}element->next = NULL;
+                ${prev}element->data = (union Data*)task;
+                ${prev}queue->last->next  = element;
+                ${prev}queue->last = element;
 EOFEOF
                 print $fd $putTask;
                 next;