Mercurial > hg > Game > Cerium

annotate example/bitonic_sort/cuda/swap.cu @ 2054:2e7a6f40672f draft

Find changesets by keywords (author, files, the commit message), revision number or hash, or revset expression.
add param(4) in FileMapReduce.cc
author masa
date Fri, 29 Jan 2016 15:56:28 +0900
parents bc010492ade4
children
Ignore whitespace changes - Everywhere: Within whitespace: At end of lines:
rev   line source
2040
bc010492ade4 add cuda task with bitonic sort.
Shohei KOKUBO <e105744@ie.u-ryukyu.ac.jp>
parents:
diff changeset
1 extern "C" {
bc010492ade4 add cuda task with bitonic sort.
Shohei KOKUBO <e105744@ie.u-ryukyu.ac.jp>
parents:
diff changeset
2 __global__ void
bc010492ade4 add cuda task with bitonic sort.
Shohei KOKUBO <e105744@ie.u-ryukyu.ac.jp>
parents:
diff changeset
3 swap(long* param, int* Data)
bc010492ade4 add cuda task with bitonic sort.
Shohei KOKUBO <e105744@ie.u-ryukyu.ac.jp>
parents:
diff changeset
4 {
bc010492ade4 add cuda task with bitonic sort.
Shohei KOKUBO <e105744@ie.u-ryukyu.ac.jp>
parents:
diff changeset
5 long block = param[0];
bc010492ade4 add cuda task with bitonic sort.
Shohei KOKUBO <e105744@ie.u-ryukyu.ac.jp>
parents:
diff changeset
6 long first = param[1];
bc010492ade4 add cuda task with bitonic sort.
Shohei KOKUBO <e105744@ie.u-ryukyu.ac.jp>
parents:
diff changeset
7 int x = blockIdx.x*blockDim.x+threadIdx.x;
bc010492ade4 add cuda task with bitonic sort.
Shohei KOKUBO <e105744@ie.u-ryukyu.ac.jp>
parents:
diff changeset
8
bc010492ade4 add cuda task with bitonic sort.
Shohei KOKUBO <e105744@ie.u-ryukyu.ac.jp>
parents:
diff changeset
9 int position = x/block;
bc010492ade4 add cuda task with bitonic sort.
Shohei KOKUBO <e105744@ie.u-ryukyu.ac.jp>
parents:
diff changeset
10 int index = x+block*position;
bc010492ade4 add cuda task with bitonic sort.
Shohei KOKUBO <e105744@ie.u-ryukyu.ac.jp>
parents:
diff changeset
11
bc010492ade4 add cuda task with bitonic sort.
Shohei KOKUBO <e105744@ie.u-ryukyu.ac.jp>
parents:
diff changeset
12 block = (first == 1) ? ((block<<1)*(position+1))-(index%block)-1 : index+block;
bc010492ade4 add cuda task with bitonic sort.
Shohei KOKUBO <e105744@ie.u-ryukyu.ac.jp>
parents:
diff changeset
13
bc010492ade4 add cuda task with bitonic sort.
Shohei KOKUBO <e105744@ie.u-ryukyu.ac.jp>
parents:
diff changeset
14 if (Data[block] < Data[index]) {
bc010492ade4 add cuda task with bitonic sort.
Shohei KOKUBO <e105744@ie.u-ryukyu.ac.jp>
parents:
diff changeset
15 int tmp = Data[index];
bc010492ade4 add cuda task with bitonic sort.
Shohei KOKUBO <e105744@ie.u-ryukyu.ac.jp>
parents:
diff changeset
16 Data[index] = Data[block];
bc010492ade4 add cuda task with bitonic sort.
Shohei KOKUBO <e105744@ie.u-ryukyu.ac.jp>
parents:
diff changeset
17 Data[block] = tmp;
bc010492ade4 add cuda task with bitonic sort.
Shohei KOKUBO <e105744@ie.u-ryukyu.ac.jp>
parents:
diff changeset
18 }
bc010492ade4 add cuda task with bitonic sort.
Shohei KOKUBO <e105744@ie.u-ryukyu.ac.jp>
parents:
diff changeset
19 }
bc010492ade4 add cuda task with bitonic sort.
Shohei KOKUBO <e105744@ie.u-ryukyu.ac.jp>
parents:
diff changeset
20 }