Mercurial > hg > Gears > GearsAgda
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; + } + } +}