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