# HG changeset patch
# User Shohei KOKUBO <e105744@ie.u-ryukyu.ac.jp>
# Date 1392103786 -32400
# Node ID 7d1afa7aeccd097947d98cf9875c5a08b25ef68f
# Parent  273de551f7262ab0f4e35bf9a266920541b76903
add file

diff -r 273de551f726 -r 7d1afa7aeccd example/many_task/Makefile.cuda
--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/example/many_task/Makefile.cuda	Tue Feb 11 16:29:46 2014 +0900
@@ -0,0 +1,55 @@
+include ./Makefile.def
+
+SRCS_TMP = $(wildcard *.cc)
+SRCS_EXCLUDE = sort_test.cc  # 除外するファイルを書く
+SRCS = $(filter-out $(SRCS_EXCLUDE),$(SRCS_TMP))
+OBJS = $(SRCS:.cc=.o)
+
+TASK_DIR  = ppe
+TASK_SRCS_TMP = $(wildcard $(TASK_DIR)/*.cc)
+TASK_SRCS_EXCLUDE = sort_test.cc
+TASK_SRCS = $(filter-out $(TASK_DIR)/$(TASK_SRCS_EXCLUDE),$(TASK_SRCS_TMP))
+TASK_OBJS = $(TASK_SRCS:.cc=.o)
+
+CUDA_TASK_DIR  = cuda
+CUDA_TASK_SRCS_TMP = $(wildcard $(CUDA_TASK_DIR)/*.cu)
+CUDA_TASK_SRCS_EXCLUDE =
+CUDA_TASK_SRCS = $(filter-out $(CUDA_TASK_DIR)/$(CUDA_TASK_SRCS_EXCLUDE),$(CUDA_TASK_SRCS_TMP))
+CUDA_TASK_OBJS = $(CUDA_TASK_SRCS:.cu=.ptx)
+
+CC += $(ABI)
+CFLAGS  += -D__CERIUM_CUDA__
+
+LIBS += -L${CERIUM}/TaskManager -DUSE_SIMPLE_TASK -lCudaManager -F/Library/Frameworks -framework CUDA `sdl-config --libs`
+
+INCLUDE += -I/Developer/NVIDIA/CUDA-5.5/include
+
+NVCC = nvcc
+NVCCFLAGS = -ptx -arch=sm_20
+
+.SUFFIXES: .cc .o .cu .ptx
+
+.cc.o:
+	$(CC) $(CFLAGS) $(INCLUDE) -c $< -o $@
+
+.cu.ptx:
+	$(NVCC) $(NVCCFLAGS) $< -o $@
+
+all: $(TARGET)
+
+$(TARGET): $(OBJS) $(TASK_OBJS) $(CUDA_TASK_OBJS)
+	$(CC) -o $@ $(OBJS) $(TASK_OBJS) $(LIBS)
+
+link:
+	$(CC) -o $(TARGET) $(OBJS) $(TASK_OBJS) $(LIBS)
+
+debug: $(TARGET)
+	sudo lldb -- ./$(TARGET) -gpu -g
+
+test : $(TARGET)
+	./$(TARGET) -gpu -g
+
+clean:
+	rm -f $(TARGET) $(OBJS) $(TASK_OBJS) $(CUDA_TASK_OBJS)
+	rm -f *~ \#*
+	rm -f cuda/*~ cuda/\#*
diff -r 273de551f726 -r 7d1afa7aeccd example/many_task/cuda/QuickSort.cu
--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/example/many_task/cuda/QuickSort.cu	Tue Feb 11 16:29:46 2014 +0900
@@ -0,0 +1,70 @@
+extern "C" {
+    typedef struct Data {
+        int index;
+        int ptr;
+        int pad[2];
+    } Data, *DataPtr;
+    
+    __device__ void
+    swap(Data *data, int left, int right )
+    {
+        Data tmp    = data[left];
+        data[left]  = data[right];
+        data[right] = tmp;
+    }
+    
+    __device__ void
+    bubble_sort(Data *data, int begin, int end)
+    {
+        for (int count=0;count<end;count++) {
+            for (int c=end;c>count;c--) {
+                if (data[c].index<data[c-1].index)swap(data,c-1,c);
+            }
+        }
+    }
+    
+    
+    __global__ void
+    quick_sort(int *count,
+               Data *data)
+    {
+        int begin = 0;
+        int end = count[0]-1;
+        
+        int stack[1024];
+        int sp = 0;
+        int p;
+        
+        while (1) {
+            while (begin < end) {
+                /*
+                 * if (end-begin <= 50) {
+                 *     bubble_sort(data,begin,end);
+                 *     break;
+                 * }
+                 */
+                int where = (begin + end) / 2;
+                int pivot = data[where].index;
+                data[where].index = data[begin].index;
+                int i;
+                p = begin;
+                for (i=begin+1; i<=end; i++) {
+                    if (data[i].index < pivot) {
+                        p++;
+                        swap(data, p, i);
+                    }
+                }
+                data[begin].index = data[p].index;
+                data[p].index = pivot;
+                
+                stack[sp++] = p + 1;
+                stack[sp++] = end;
+                end = p - 1;
+            }
+            if (sp == 0) return;
+            end = stack[--sp];
+            begin = stack[--sp];
+            begin = p + 1;
+        }
+    }
+}