Mercurial > hg > Game > Cerium
changeset 1972:6fa9e5d55774 draft
add file
| author | Shohei KOKUBO <e105744@ie.u-ryukyu.ac.jp> |
|---|---|
| date | Wed, 26 Feb 2014 13:10:20 +0900 |
| parents | fdb3ed0bc51d |
| children | 6dd11261489a |
| files | example/fft/cuda/bitReverse.cu example/fft/cuda/task_init.cc |
| diffstat | 2 files changed, 43 insertions(+), 0 deletions(-) [+] |
line wrap: on
line diff
--- /dev/null Thu Jan 01 00:00:00 1970 +0000 +++ b/example/fft/cuda/bitReverse.cu Wed Feb 26 13:10:20 2014 +0900 @@ -0,0 +1,29 @@ +extern "C" { +#ifdef __APPLE__ +#include <OpenCL/opencl.h> +#else +#include <CL/cl.h> +#endif + + __global__ void + bitReverse(long* param, cl_float2* src, cl_float2* dst) + { + unsigned long gid = blockIdx.x*blockDim.x+threadIdx.x; // (unsigned long)s->get_param(0); + unsigned long nid = blockIdx.y*blockDim.y+threadIdx.y; // (unsigned long)s->get_param(1); + + unsigned int j = gid; + + unsigned long m = param[0]; + unsigned long n = param[1]; + + j = (j & 0x55555555) << 1 | (j & 0xAAAAAAAA) >> 1; + j = (j & 0x33333333) << 2 | (j & 0xCCCCCCCC) >> 2; + j = (j & 0x0F0F0F0F) << 4 | (j & 0xF0F0F0F0) >> 4; + j = (j & 0x00FF00FF) << 8 | (j & 0xFF00FF00) >> 8; + j = (j & 0x0000FFFF) << 16 | (j & 0xFFFF0000) >> 16; + + j >>= (32-m); + + dst[nid*n+j] = src[nid*n+gid]; + } +}
--- /dev/null Thu Jan 01 00:00:00 1970 +0000 +++ b/example/fft/cuda/task_init.cc Wed Feb 26 13:10:20 2014 +0900 @@ -0,0 +1,14 @@ +#include "Func.h" +#include "Scheduler.h" +#include "CudaScheduler.h" + +void +gpu_task_init(void) +{ + CudaSchedRegisterTask(SPIN_FACT, "cuda/spinFact.ptx", "spinFact"); + CudaSchedRegisterTask(NORMALIZATION, "cuda/norm.ptx", "norm"); + CudaSchedRegisterTask(BIT_REVERSE, "cuda/bitReverse.ptx", "bitReverse"); + CudaSchedRegisterTask(BUTTERFLY, "cuda/butterfly.ptx", "butterfly"); + CudaSchedRegisterTask(TRANSPOSE, "cuda/transpose.ptx", "transpose"); + CudaSchedRegisterTask(HIGH_PASS_FILTER, "gpu/highPassFilter.ptx", "highPassFilter"); +}