# HG changeset patch # User Tatsuki IHA # Date 1492463062 -32400 # Node ID 8592a1d161b7217921418175a692393caa48e268 # Parent 48c2b5bcab79929a6a049cbee67c9189740508bf Move CUDAtwice example to examples directory diff -r 48c2b5bcab79 -r 8592a1d161b7 src/parallel_execution/CMakeLists.txt --- a/src/parallel_execution/CMakeLists.txt Tue Apr 18 06:00:45 2017 +0900 +++ b/src/parallel_execution/CMakeLists.txt Tue Apr 18 06:04:22 2017 +0900 @@ -75,7 +75,7 @@ TARGET CUDAtwice SOURCES - main.cbc RedBlackTree.cbc compare.c SingleLinkedStack.cbc CPUWorker.cbc time.cbc twice.cbc TaskManagerImpl.cbc SingleLinkedQueue.cbc SynchronizedQueue.cbc SemaphoreImpl.cbc CUDAWorker.cbc CUDAtwice.cbc CUDAtwice.cu cuda.c + main.cbc RedBlackTree.cbc compare.c SingleLinkedStack.cbc CPUWorker.cbc time.cbc twice.cbc TaskManagerImpl.cbc SingleLinkedQueue.cbc SynchronizedQueue.cbc SemaphoreImpl.cbc CUDAWorker.cbc examples/CUDAtwice.cbc examples/CUDAtwice.cu cuda.c ) set_target_properties(CUDAtwice PROPERTIES COMPILE_FLAGS "-Wall -g -DUSE_CUDAWorker=1 -DUSE_CUDA_MAIN_THREAD") endif() diff -r 48c2b5bcab79 -r 8592a1d161b7 src/parallel_execution/CUDAtwice.cbc --- a/src/parallel_execution/CUDAtwice.cbc Tue Apr 18 06:00:45 2017 +0900 +++ /dev/null Thu Jan 01 00:00:00 1970 +0000 @@ -1,28 +0,0 @@ -#include -#include "../context.h" - - -extern void CUDAExec(struct Context* context, Array* array, LoopCounter *loopCounter); - -__code CUDAtwice(struct Context* context, struct LoopCounter* loopCounter, int index, int prefix, int* array, struct Context* workerContext) { - int i = loopCounter->i; - if (i < prefix) { - array[i+index*prefix] = array[i+index*prefix]*2; - loopCounter->i++; - - goto meta(context, C_twice); - } - - loopCounter->i = 0; - goto meta(workerContext, workerContext->next); -} - -__code CUDAtwice_stub(struct Context* context) { -printf("CUdAtwice stub\n"); - struct LoopCounter* loopCounter = &context->data[context->dataNum]->LoopCounter; - struct Array* array = &context->data[context->dataNum+1]->Array; - CUDAExec(context,array,loopCounter); - - //continuationにそってGPUworkerに戻る - goto meta(context, context->next); -} diff -r 48c2b5bcab79 -r 8592a1d161b7 src/parallel_execution/CUDAtwice.cu --- a/src/parallel_execution/CUDAtwice.cu Tue Apr 18 06:00:45 2017 +0900 +++ /dev/null Thu Jan 01 00:00:00 1970 +0000 @@ -1,33 +0,0 @@ -extern "C" { - -#include - -// __global__ void twice(struct LoopCounter* loopCounter, int prefix ,int* array) { -// int index = blockIdx.x * blockDim.x + threadIdx.x; -// printf("array %p, blockIdx.x = %d, blockDim.x = %d, threadIdx.x = %d\n"); -// int i = 0; -// while (i < prefix) { -// array[i+index*prefix] = array[i+index*prefix]*2; -// } -// } - - struct LoopCounter { - int i; - } LoopCounter; - - __global__ void twice(struct LoopCounter* loopCounter, int index, int prefix, int* array) { - printf("array %p, index = %d, prefix = %d loopCounter->i %d\n",array,index,prefix,loopCounter->i); -C_twice: - int i = loopCounter->i; - if (i < prefix) { - array[i+index*prefix] = array[i+index*prefix]*2; - loopCounter->i++; - - goto C_twice; - } - - loopCounter->i = 0; - } - - -} diff -r 48c2b5bcab79 -r 8592a1d161b7 src/parallel_execution/examples/CUDAtwice.cbc --- /dev/null Thu Jan 01 00:00:00 1970 +0000 +++ b/src/parallel_execution/examples/CUDAtwice.cbc Tue Apr 18 06:04:22 2017 +0900 @@ -0,0 +1,28 @@ +#include +#include "../../context.h" + + +extern void CUDAExec(struct Context* context, Array* array, LoopCounter *loopCounter); + +__code CUDAtwice(struct Context* context, struct LoopCounter* loopCounter, int index, int prefix, int* array, struct Context* workerContext) { + int i = loopCounter->i; + if (i < prefix) { + array[i+index*prefix] = array[i+index*prefix]*2; + loopCounter->i++; + + goto meta(context, C_twice); + } + + loopCounter->i = 0; + goto meta(workerContext, workerContext->next); +} + +__code CUDAtwice_stub(struct Context* context) { +printf("CUdAtwice stub\n"); + struct LoopCounter* loopCounter = &context->data[context->dataNum]->LoopCounter; + struct Array* array = &context->data[context->dataNum+1]->Array; + CUDAExec(context,array,loopCounter); + + //continuationにそってGPUworkerに戻る + goto meta(context, context->next); +} diff -r 48c2b5bcab79 -r 8592a1d161b7 src/parallel_execution/examples/CUDAtwice.cu --- /dev/null Thu Jan 01 00:00:00 1970 +0000 +++ b/src/parallel_execution/examples/CUDAtwice.cu Tue Apr 18 06:04:22 2017 +0900 @@ -0,0 +1,33 @@ +extern "C" { + +#include + +// __global__ void twice(struct LoopCounter* loopCounter, int prefix ,int* array) { +// int index = blockIdx.x * blockDim.x + threadIdx.x; +// printf("array %p, blockIdx.x = %d, blockDim.x = %d, threadIdx.x = %d\n"); +// int i = 0; +// while (i < prefix) { +// array[i+index*prefix] = array[i+index*prefix]*2; +// } +// } + + struct LoopCounter { + int i; + } LoopCounter; + + __global__ void twice(struct LoopCounter* loopCounter, int index, int prefix, int* array) { + printf("array %p, index = %d, prefix = %d loopCounter->i %d\n",array,index,prefix,loopCounter->i); +C_twice: + int i = loopCounter->i; + if (i < prefix) { + array[i+index*prefix] = array[i+index*prefix]*2; + loopCounter->i++; + + goto C_twice; + } + + loopCounter->i = 0; + } + + +}