Mercurial > hg > Game > Cerium

changeset 1665:bba05acb7a16 draft

Find changesets by keywords (author, files, the commit message), revision number or hash, or revset expression.
merge
author Masa <e085726@ie.u-ryukyu.ac.jp>
date Tue, 16 Jul 2013 20:16:43 +0900
parents 99f8130793b4 (current diff) ce031df3dd32 (diff)
children 805f60a3e9f9
files
diffstat 4 files changed, 33 insertions(+), 35 deletions(-) [+]
line wrap: on
line diff
--- a/TaskManager/test/GpuRunTest/GpuRunTest.cc	Tue Jul 16 20:14:10 2013 +0900
+++ b/TaskManager/test/GpuRunTest/GpuRunTest.cc	Tue Jul 16 20:16:43 2013 +0900
@@ -8,8 +8,8 @@
 #include "Func.h"
 
 #define DEFAULT 5
-static long length = DEFAULT;
-static long times = 2;
+static int length = DEFAULT;
+static int times = 2;
 static int task = 1;
 
 extern void gpu_task_init(void);
@@ -86,7 +86,7 @@
 
     HTaskPtr twice = manager->create_task(Twice);
     twice->set_param(0, (memaddr)length);
-    twice->set_param(1, (memaddr)times);
+    twice->set_param(3, (memaddr)times);
     /*
     twice->set_param(2, (memaddr)2);
     twice->set_param(3, (memaddr)3);
--- a/TaskManager/test/GpuRunTest/gpu/twice.cl	Tue Jul 16 20:14:10 2013 +0900
+++ b/TaskManager/test/GpuRunTest/gpu/twice.cl	Tue Jul 16 20:16:43 2013 +0900
@@ -4,7 +4,7 @@
       __global int *output_data)
 {
     long count = (long)param[0];
-    long times = (long)param[1];
+    long times = (long)param[3];
     for (int i = 0; i<count; i++) {
         output_data[i] = input_data[i] * times; 
     }
--- a/example/fft/gpu/fft.cl	Tue Jul 16 20:14:10 2013 +0900
+++ b/example/fft/gpu/fft.cl	Tue Jul 16 20:16:43 2013 +0900
@@ -1,19 +1,18 @@
 #define PI 3.14159265358979323846
 #define PI_2 1.57079632679489661923
 
-__kernel void spinFact(__constant int *data_count,__constant int *n_, __global float2* w)
+__kernel void spinFact(__constant long *param,__global float2* w)
 {
-    int n = n_[0];
+    long n = param[3];
     unsigned int i = get_global_id(0);
 
     float2 angle = (float2)(2*i*PI/(float)n,(2*i*PI/(float)n)+PI_2);
     w[i] = cos(angle);
 }
 
-__kernel void bitReverse(__constant int *data_count,__global float2 *dst, __global float2 *src,__constant int *m_, __constant int *n_)
-{
-    int m = m_[0];
-    int n = n_[0];
+__kernel void bitReverse(__constant int *param, __global float2 *dst, __global float2 *src) {
+    int m = param[3];
+    int n = param[4];
     unsigned int gid = get_global_id(0);
     unsigned int nid = get_global_id(1);
 
@@ -29,23 +28,22 @@
     dst[nid*n+j] = src[nid*n+gid];
 }
 
-__kernel void norm(__constant int *data_count, __global float2 *x, __constant int *n_)
+__kernel void norm(__constant int *param, __global float2 *x,__global float2 *x_out)
 {
-    int n = n_[0];
+    int n = param[3];
     unsigned int gid = get_global_id(0);
 
 
     unsigned int nid = get_global_id(1);
 
-    x[nid*n+gid] = x[nid*n+gid] / (float2)((float)n, (float)n);
+    x_out[nid*n+gid] = x[nid*n+gid] / (float2)((float)n, (float)n);
 }
 
-__kernel void butterfly(__constant int *data_count,__global float2 *x, __global float2* w,__constant int* m_,__constant int *n_,__constant int *iter_,__constant uint *flag_)
+__kernel void butterfly(__constant int *param, __global float2 *x, __global float2* w,__global float2 *x_out)
 {
-    int m = m_[0];
-    int n = n_[0];
-    int iter = iter_[0];
-    uint flag =flag_[0];
+    int n     = param[3];
+    uint flag = param[4];
+    int iter  = param[5];
     
     unsigned int gid = get_global_id(0);
     unsigned int nid = get_global_id(1);
@@ -75,13 +73,13 @@
     resa = xa + xbxx*wab + xbyy*wayx;
     resb = xa - xbxx*wab + xbyy*wbyx;
 
-    x[a] = resa;
-    x[b] = resb;
+    x_out[a] = resa;
+    x_out[b] = resb;
 }
 
-__kernel void transpose(__constant int *data_count, __global float2* src,__constant int *n_, __global float2 *dst)
+__kernel void transpose(__constant int *param, __global float2* src,__global float2 *dst)
 {
-    int n = n_[0];
+    int n = param[3];
     unsigned int xgid = get_global_id(0);
     unsigned int ygid = get_global_id(1);
 
@@ -91,10 +89,10 @@
     dst[oid] = src[iid];
 }
 
-__kernel void highPassFilter(__constant int *data_count,__constant int *n_,__constant int *radius_,__global float2* image)
+__kernel void highPassFilter(__constant int *param, __global float2* image, __global float2* image_out)
 {
-    int n = n_[0];
-    int radius = radius_[0];
+    int n = param[3];
+    int radius = param[4];
     unsigned int xgid = get_global_id(0);
     unsigned int ygid = get_global_id(1);
 
@@ -114,5 +112,5 @@
     } else {
         window = (int2)(-1L, -1L);
     }
-        image[ygid*n+xgid] = as_float2(as_int2(image[ygid*n+xgid]) & window);
+        image_out[ygid*n+xgid] = as_float2(as_int2(image[ygid*n+xgid]) & window);
 }
--- a/example/fft/main.cc	Tue Jul 16 20:14:10 2013 +0900
+++ b/example/fft/main.cc	Tue Jul 16 20:16:43 2013 +0900
@@ -105,24 +105,24 @@
 }
 
 HTask*
-fftCore(TaskManager *manager,cl_float2 *dst, cl_float2 *src, cl_float2 *spin, int m, enum Mode direction,HTask* waitTask)
+fftCore(TaskManager *manager,cl_float2 *dst, cl_float2 *src, cl_float2 *spin, long m, enum Mode direction,HTask* waitTask)
 {
     unsigned int direction_flag;
     switch (direction) {
     case forward:direction_flag = 0x00000000; break;
     case inverse:direction_flag = 0x80000000; break;
     }
-    int n = 1<<m;
+    long n = 1<<m;
     size_t gws[2],lws[2];
     int length_dst = n*n;
     int length_src = n*n;
 
     HTask* brev = manager->create_task(BIT_REVERSE);
     setWorkSize(gws,lws,n,n);
+    brev->set_param(3,m);
+    brev->set_param(4,n);
     brev->set_inData(0, src, length_src*sizeof(cl_float2));
     brev->set_outData(0, dst, length_dst*sizeof(cl_float2));
-    brev->set_param(3,m);
-    brev->set_param(4,n);
     brev->set_cpu(spe_cpu);
     brev->wait_for(waitTask);
     brev->iterate(gws[0],gws[1]);
@@ -132,12 +132,12 @@
     setWorkSize(gws,lws,n/2,n);
     for(int iter=1;iter<=m;iter++) {
         HTask* bfly = manager->create_task(BUTTERFLY);
+        bfly->set_param(3,n);
+        bfly->set_param(4,direction_flag);
+        bfly->set_param(5,(long)iter);
         bfly->set_inData(0, dst, length_dst*sizeof(cl_float2));
         bfly->set_inData(1, spin, sizeof(cl_float2)*(n/2));
         bfly->set_outData(0,dst,length_dst*sizeof(cl_float2));
-        bfly->set_param(3,n);
-        bfly->set_param(4,direction_flag);
-        bfly->set_param(5,iter);
         bfly->set_cpu(spe_cpu);
         bfly->wait_for(waitTask);
         bfly->iterate(gws[0],gws[1]);
@@ -182,8 +182,8 @@
 void
 run_start(TaskManager *manager,pgm_t ipgm)
 {
-    int n = ipgm.width;
-    int m = (cl_int)(log((double)n)/log(2.0));
+    long n = ipgm.width;
+    long m = (cl_int)(log((double)n)/log(2.0));
     size_t *gws = new size_t[2];
     size_t *lws = new size_t[2];