Mercurial > hg > Gears > Gears
changeset 404:c5cd9888bf2a
Fix bitonicSort
| author | Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> |
|---|---|
| date | Sun, 03 Sep 2017 00:21:16 +0900 |
| parents | 83c9aeb1fe3e |
| children | 8915fce522b3 |
| files | src/parallel_execution/CMakeLists.txt src/parallel_execution/context.h src/parallel_execution/cuda.c src/parallel_execution/examples/bitonicSort/bitonicSort.cbc src/parallel_execution/examples/bitonicSort/bitonicSwap.cbc src/parallel_execution/examples/bitonicSort/makeArray.cbc src/parallel_execution/examples/bitonicSort/printArray.cbc src/parallel_execution/examples/bitonicSort/swap.cbc src/parallel_execution/generate_stub.pl |
| diffstat | 9 files changed, 147 insertions(+), 195 deletions(-) [+] |
line wrap: on
line diff
--- 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})
--- 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;
--- 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)); //結果を取ってくるコマンドを入力する //コマンドの終了待ちを行う
--- 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); -}
--- /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)); +}
--- /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); +}
--- /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); +}
--- 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); -}
--- 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;