From b3bea947efa55d2c0f198b6c6795b3177be27f45 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Picca=20Fr=C3=A9d=C3=A9ric-Emmanuel?= Date: Wed, 6 Jan 2021 14:10:12 +0100 Subject: New upstream version 0.14.0+dfsg --- silx/resources/gui/icons/add.png | Bin 0 -> 470 bytes silx/resources/gui/icons/add.svg | 2 + silx/resources/gui/icons/backend-opengl.png | Bin 0 -> 1582 bytes silx/resources/gui/icons/backend-opengl.svg | 18 ++ silx/resources/gui/icons/rm.png | Bin 0 -> 348 bytes silx/resources/gui/icons/rm.svg | 2 + silx/resources/opencl/backproj.cl | 301 +++------------------------- silx/resources/opencl/proj.cl | 4 +- 8 files changed, 48 insertions(+), 279 deletions(-) create mode 100644 silx/resources/gui/icons/add.png create mode 100644 silx/resources/gui/icons/add.svg create mode 100644 silx/resources/gui/icons/backend-opengl.png create mode 100644 silx/resources/gui/icons/backend-opengl.svg create mode 100644 silx/resources/gui/icons/rm.png create mode 100644 silx/resources/gui/icons/rm.svg (limited to 'silx/resources') diff --git a/silx/resources/gui/icons/add.png b/silx/resources/gui/icons/add.png new file mode 100644 index 0000000..80c6400 Binary files /dev/null and b/silx/resources/gui/icons/add.png differ diff --git a/silx/resources/gui/icons/add.svg b/silx/resources/gui/icons/add.svg new file mode 100644 index 0000000..19c1a6d --- /dev/null +++ b/silx/resources/gui/icons/add.svg @@ -0,0 +1,2 @@ + +image/svg+xml diff --git a/silx/resources/gui/icons/backend-opengl.png b/silx/resources/gui/icons/backend-opengl.png new file mode 100644 index 0000000..ff81f64 Binary files /dev/null and b/silx/resources/gui/icons/backend-opengl.png differ diff --git a/silx/resources/gui/icons/backend-opengl.svg b/silx/resources/gui/icons/backend-opengl.svg new file mode 100644 index 0000000..41d79b8 --- /dev/null +++ b/silx/resources/gui/icons/backend-opengl.svg @@ -0,0 +1,18 @@ + + + + + + image/svg+xml + + + + + + + + + + + + diff --git a/silx/resources/gui/icons/rm.png b/silx/resources/gui/icons/rm.png new file mode 100644 index 0000000..ecff08b Binary files /dev/null and b/silx/resources/gui/icons/rm.png differ diff --git a/silx/resources/gui/icons/rm.svg b/silx/resources/gui/icons/rm.svg new file mode 100644 index 0000000..7cc515e --- /dev/null +++ b/silx/resources/gui/icons/rm.svg @@ -0,0 +1,2 @@ + +image/svg+xml diff --git a/silx/resources/opencl/backproj.cl b/silx/resources/opencl/backproj.cl index 6fadc2c..da15131 100644 --- a/silx/resources/opencl/backproj.cl +++ b/silx/resources/opencl/backproj.cl @@ -35,7 +35,7 @@ /************************ GPU VERSION (with textures) **************************/ /*******************************************************************************/ - +#ifndef DONT_USE_TEXTURES kernel void backproj_kernel( int num_proj, int num_bins, @@ -55,11 +55,6 @@ kernel void backproj_kernel( const int tidy = get_local_id(1); //threadIdx.y; const int bidy = get_group_id(1); //blockIdx.y; - //~ local float shared[768]; - //~ float * sh_sin = shared; - //~ float * sh_cos = shared+256; - //~ float * sh_axis = sh_cos+256; - local float sh_cos[256]; local float sh_sin[256]; local float sh_axis[256]; @@ -107,7 +102,7 @@ kernel void backproj_kernel( d_SLICE[ 32*get_num_groups(0)*(bidy*32+tidy*2+0) + bidx*32 + tidx*2 + 1] = res2; d_SLICE[ 32*get_num_groups(0)*(bidy*32+tidy*2+1) + bidx*32 + tidx*2 + 1] = res3; } - +#endif @@ -134,7 +129,7 @@ static float linear_interpolation(float2 vals, { if (xm == xp) return vals.s0; - else + else return (vals.s0 * (xp - x)) + (vals.s1 * (x - xm)); } @@ -197,280 +192,36 @@ kernel void backproj_cpu_kernel( h1 = (acorr05 + (bx00+0)*pcos - (by00+1)*psin); h2 = (acorr05 + (bx00+1)*pcos - (by00+0)*psin); h3 = (acorr05 + (bx00+1)*pcos - (by00+1)*psin); - - - float x; - int ym, xm, xp; - ym = proj; - float2 vals; - - if(h0>=0 && h0=0 && h1=0 && h2=0 && h3= Nx) adj_coords.s1 = Nx - 1; - if (adj_coords.s2 < 0) adj_coords.s2 = 0; - if (adj_coords.s3 >= Ny) adj_coords.s3 = Ny -1; - if (adj_coords.s0 >= Nx) adj_coords.s0 = Nx - 1; - if (adj_coords.s2 >= Ny) adj_coords.s2 = Ny -1; - // Interp - val = adj_vals.s1*(adj_coords.s1-x)*(y-adj_coords.s2) - + adj_vals.s2 *(x-adj_coords.s0)*(y-adj_coords.s2) - + adj_vals.s0 *(adj_coords.s1-x)*(adj_coords.s3-y) - + adj_vals.s3 *(x-adj_coords.s0)*(adj_coords.s3-y); - - } - return val; -} -*/ - - -/* -__kernel void backproj_cpu_kernel_good( - int num_proj, - int num_bins, - float axis_position, - __global float *d_SLICE, - __global float* d_sino, - float gpu_offset_x, - float gpu_offset_y, - __global float * d_cos_s, // precalculated cos(theta[i]) - __global float * d_sin_s, // precalculated sin(theta[i]) - __global float * d_axis_s, // array of axis positions (n_projs) - __local float* shared2) // 768B of local mem -{ - const int tidx = get_local_id(0); //threadIdx.x; - const int bidx = get_group_id(0); //blockIdx.x; - const int tidy = get_local_id(1); //threadIdx.y; - const int bidy = get_group_id(1); //blockIdx.y; - - //~ __local float shared[768]; - //~ float * sh_sin = shared; - //~ float * sh_cos = shared+256; - //~ float * sh_axis = sh_cos+256; - __local float sh_cos[256]; - __local float sh_sin[256]; - __local float sh_axis[256]; - - float pcos, psin; - float h0, h1, h2, h3; - const float apos_off_x= gpu_offset_x - axis_position ; - const float apos_off_y= gpu_offset_y - axis_position ; - float acorr05; - float res0 = 0, res1 = 0, res2 = 0, res3 = 0; - const float bx00 = (32 * bidx + 2 * tidx + 0 + apos_off_x ) ; - const float by00 = (32 * bidy + 2 * tidy + 0 + apos_off_y ) ; + float x; + int ym, xm, xp; + ym = proj; + float2 vals; - int read=0; - for(int proj=0; proj=read) { - barrier(CLK_LOCAL_MEM_FENCE); - int ip = tidy*16+tidx; - if( read+ip < num_proj) { - sh_cos [ip] = d_cos_s[read+ip] ; - sh_sin [ip] = d_sin_s[read+ip] ; - sh_axis[ip] = d_axis_s[read+ip] ; - } - read=read+256; // 256=16*16 block size - barrier(CLK_LOCAL_MEM_FENCE); + if(h0>=0 && h0=0 && h0=0 && h1=0 && h2=0 && h3