diff src/parallel_execution/examples/bitonicSort/CUDAbitonicSwap.cu @ 414:49159fbdd1fb

Work CUDAbitonicSort
author Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
date Fri, 15 Sep 2017 22:49:45 +0900
parents
children 764c92c3b181
line wrap: on
line diff
--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/src/parallel_execution/examples/bitonicSort/CUDAbitonicSwap.cu	Fri Sep 15 22:49:45 2017 +0900
@@ -0,0 +1,25 @@
+extern "C" {
+    struct Integer {
+        int value;
+    };
+    __global__ void bitonicSwap(struct Integer* array, int* blockPtr, int* firstPtr, int* prefixPtr) {
+        int block = *blockPtr;
+        int first = *firstPtr;
+        int prefix = *prefixPtr;
+        int i = 0;
+C_bitonicSwap:
+        if (i < prefix) {
+            int index = i + blockIdx.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;
+        }
+    }
+}