annotate src/parallel_execution/examples/bitonicSort/CUDAbitonicSwap.cu @ 420:764c92c3b181

Fix
author Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
date Fri, 06 Oct 2017 15:10:07 +0900 (2017-10-06)
parents 49159fbdd1fb
children 6bb391fc9e12
Ignore whitespace changes - Everywhere: Within whitespace: At end of lines:
rev   line source
414
49159fbdd1fb Work CUDAbitonicSort
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
1 extern "C" {
49159fbdd1fb Work CUDAbitonicSort
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
2 struct Integer {
49159fbdd1fb Work CUDAbitonicSort
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
3 int value;
49159fbdd1fb Work CUDAbitonicSort
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
4 };
49159fbdd1fb Work CUDAbitonicSort
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
5 __global__ void bitonicSwap(struct Integer* array, int* blockPtr, int* firstPtr, int* prefixPtr) {
49159fbdd1fb Work CUDAbitonicSort
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
6 int block = *blockPtr;
49159fbdd1fb Work CUDAbitonicSort
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
7 int first = *firstPtr;
49159fbdd1fb Work CUDAbitonicSort
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
8 int prefix = *prefixPtr;
49159fbdd1fb Work CUDAbitonicSort
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
9 int i = 0;
49159fbdd1fb Work CUDAbitonicSort
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
10 C_bitonicSwap:
49159fbdd1fb Work CUDAbitonicSort
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
11 if (i < prefix) {
420
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents: 414
diff changeset
12 int index = i + (blockIdx.x * blockDim.x + threadIdx.x) * prefix;
414
49159fbdd1fb Work CUDAbitonicSort
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
13 int position = index/block;
49159fbdd1fb Work CUDAbitonicSort
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
14 int index1 = index+block*position;
49159fbdd1fb Work CUDAbitonicSort
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
15 int index2 = (first == 1)? ((block<<1)*(position+1))-(index1%block)-1 : index1+block;
49159fbdd1fb Work CUDAbitonicSort
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
16 if (array[index2].value < array[index1].value) {
49159fbdd1fb Work CUDAbitonicSort
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
17 struct Integer tmp = array[index1];
49159fbdd1fb Work CUDAbitonicSort
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
18 array[index1] = array[index2];
49159fbdd1fb Work CUDAbitonicSort
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
19 array[index2] = tmp;
49159fbdd1fb Work CUDAbitonicSort
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
20 }
49159fbdd1fb Work CUDAbitonicSort
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
21 i++;
49159fbdd1fb Work CUDAbitonicSort
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
22 goto C_bitonicSwap;
49159fbdd1fb Work CUDAbitonicSort
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
23 }
49159fbdd1fb Work CUDAbitonicSort
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
24 }
49159fbdd1fb Work CUDAbitonicSort
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
25 }