Mercurial > hg > Game > Cerium

annotate example/fft/gpu/bitReverse.cl @ 1876:5e17ab506299 draft

Find changesets by keywords (author, files, the commit message), revision number or hash, or revset expression.
change CPU_TYPE SPE_ANY to IO/0 ( cannot running )
author Masataka Kohagura <e085726@ie.u-ryukyu.ac.jp>
date Mon, 30 Dec 2013 20:52:21 +0900
parents 144e573b030b
children
Ignore whitespace changes - Everywhere: Within whitespace: At end of lines:
rev   line source
1673
2c3adce7eb40 fix fft on gpu
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
1 __kernel
2c3adce7eb40 fix fft on gpu
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
2 void bitReverse(__constant long *param, __global float2 *src,__global float2 *dst)
2c3adce7eb40 fix fft on gpu
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
3 {
2c3adce7eb40 fix fft on gpu
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
4 unsigned long gid = (unsigned long)get_global_id(0);
2c3adce7eb40 fix fft on gpu
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
5 unsigned long nid = (unsigned long)get_global_id(1);
2c3adce7eb40 fix fft on gpu
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
6
1835
144e573b030b fix fft
Shohei KOKUBO <e105744@ie.u-ryukyu.ac.jp>
parents: 1673
diff changeset
7 unsigned long m = (unsigned long)param[0];
144e573b030b fix fft
Shohei KOKUBO <e105744@ie.u-ryukyu.ac.jp>
parents: 1673
diff changeset
8 unsigned long n = (unsigned long)param[1];
1673
2c3adce7eb40 fix fft on gpu
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
9 unsigned int j = gid;
2c3adce7eb40 fix fft on gpu
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
10
2c3adce7eb40 fix fft on gpu
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
11 j = (j & 0x55555555) << 1 | (j & 0xAAAAAAAA) >> 1;
2c3adce7eb40 fix fft on gpu
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
12 j = (j & 0x33333333) << 2 | (j & 0xCCCCCCCC) >> 2;
2c3adce7eb40 fix fft on gpu
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
13 j = (j & 0x0F0F0F0F) << 4 | (j & 0xF0F0F0F0) >> 4;
2c3adce7eb40 fix fft on gpu
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
14 j = (j & 0x00FF00FF) << 8 | (j & 0xFF00FF00) >> 8;
2c3adce7eb40 fix fft on gpu
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
15 j = (j & 0x0000FFFF) << 16 | (j & 0xFFFF0000) >> 16;
2c3adce7eb40 fix fft on gpu
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
16
2c3adce7eb40 fix fft on gpu
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
17 j >>= (32-m);
2c3adce7eb40 fix fft on gpu
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
18
2c3adce7eb40 fix fft on gpu
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
19 dst[nid*n+j] = src[nid*n+gid];
2c3adce7eb40 fix fft on gpu
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
20 }