summaryrefslogtreecommitdiff
path: root/silx/resources
diff options
context:
space:
mode:
Diffstat (limited to 'silx/resources')
-rw-r--r--silx/resources/gui/icons/add.pngbin0 -> 470 bytes
-rw-r--r--silx/resources/gui/icons/add.svg2
-rw-r--r--silx/resources/gui/icons/backend-opengl.pngbin0 -> 1582 bytes
-rw-r--r--silx/resources/gui/icons/backend-opengl.svg18
-rw-r--r--silx/resources/gui/icons/rm.pngbin0 -> 348 bytes
-rw-r--r--silx/resources/gui/icons/rm.svg2
-rw-r--r--silx/resources/opencl/backproj.cl301
-rw-r--r--silx/resources/opencl/proj.cl4
8 files changed, 48 insertions, 279 deletions
diff --git a/silx/resources/gui/icons/add.png b/silx/resources/gui/icons/add.png
new file mode 100644
index 0000000..80c6400
--- /dev/null
+++ b/silx/resources/gui/icons/add.png
Binary files 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 @@
+<?xml version="1.0" encoding="UTF-8"?>
+<svg id="svg2" enable-background="new 0 0 32 32" version="1.1" viewBox="0 0 32 32" xml:space="preserve" xmlns="http://www.w3.org/2000/svg" xmlns:cc="http://creativecommons.org/ns#" xmlns:dc="http://purl.org/dc/elements/1.1/" xmlns:rdf="http://www.w3.org/1999/02/22-rdf-syntax-ns#"><metadata id="metadata24"><rdf:RDF><cc:Work rdf:about=""><dc:format>image/svg+xml</dc:format><dc:type rdf:resource="http://purl.org/dc/dcmitype/StillImage"/><dc:title/></cc:Work></rdf:RDF></metadata><rect id="rect3059" x="4.2759" y="4.2462" width="23.381" height="23.236" color="#000000" enable-background="accumulate" fill="#fff" stroke="#000" stroke-miterlimit="10" stroke-width="1.912"/><path id="path3877" d="m7.9205 16.397h15.555" color="#000000" enable-background="accumulate" fill="#fff" stroke="#000" stroke-miterlimit="10" stroke-width="3.9434"/><path id="path3877-5" d="m15.697 8.6682v15.459" color="#000000" enable-background="accumulate" fill="#fff" stroke="#000" stroke-miterlimit="10" stroke-width="3.9434"/></svg>
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
--- /dev/null
+++ b/silx/resources/gui/icons/backend-opengl.png
Binary files 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 @@
+<?xml version="1.0" encoding="UTF-8"?>
+<svg id="svg14" version="1.1" viewBox="0 0 32 32" xmlns="http://www.w3.org/2000/svg" xmlns:cc="http://creativecommons.org/ns#" xmlns:dc="http://purl.org/dc/elements/1.1/" xmlns:rdf="http://www.w3.org/1999/02/22-rdf-syntax-ns#">
+ <metadata id="metadata20">
+ <rdf:RDF>
+ <cc:Work rdf:about="">
+ <dc:format>image/svg+xml</dc:format>
+ <dc:type rdf:resource="http://purl.org/dc/dcmitype/StillImage"/>
+ </cc:Work>
+ </rdf:RDF>
+ </metadata>
+ <g id="g843" transform="matrix(1.1081 0 0 1.1081 -.938 -1.4433)" fill="#4fa6db">
+ <path id="path14" d="m7.2286 17.419c0.091251 0.56906 0.2345 1.0811 0.43 1.5362 0.19532 0.45525 0.44605 0.82097 0.7523 1.0976 0.306 0.27643 0.67067 0.41427 1.0942 0.41427 0.42335 0 0.78803-0.13784 1.0943-0.41427 0.306-0.27662 0.55673-0.64233 0.75226-1.0976 0.19528-0.45507 0.33853-0.96713 0.42978-1.5362 0.09119-0.56913 0.13682-1.1463 0.13682-1.7315s-0.04564-1.1619-0.13682-1.731c-0.09125-0.56913-0.2345-1.0811-0.42978-1.5366-0.19554-0.45494-0.44627-0.82066-0.75226-1.0972-0.30625-0.27649-0.67092-0.41477-1.0943-0.41477-0.42357 0-0.78825 0.13828-1.0942 0.41477-0.30625 0.27655-0.55698 0.64227-0.7523 1.0972-0.1955 0.45544-0.33875 0.96744-0.43 1.5366-0.091212 0.56913-0.1366 1.1458-0.1366 1.731s0.045389 1.1623 0.1366 1.7315zm-0.85106-4.0175c0.13984-0.73075 0.34979-1.3779 0.62942-1.9409 0.27963-0.56298 0.62782-1.0126 1.0441-1.3481 0.41627-0.33554 0.90103-0.50368 1.4541-0.50368 0.55284 0 1.0376 0.16814 1.4539 0.50368 0.41624 0.33542 0.76443 0.78508 1.0441 1.3481 0.27963 0.56298 0.48959 1.2102 0.62938 1.9409 0.13984 0.73075 0.20995 1.4914 0.20995 2.2818 0 0.79042-0.07012 1.5514-0.20995 2.282-0.1398 0.73075-0.34975 1.378-0.62938 1.941-0.27963 0.56298-0.62782 1.0098-1.0441 1.34-0.41628 0.32983-0.90107 0.49524-1.4539 0.49524-0.55309 0-1.0379-0.16541-1.4541-0.49524-0.41624-0.33014-0.76443-0.77701-1.0441-1.34-0.27964-0.56304-0.48959-1.2103-0.62942-1.941-0.1398-0.73063-0.20974-1.4916-0.20974-2.282 0-0.79048 0.06994-1.5511 0.20974-2.2818"/>
+ <path id="path22" d="m18.959 22.505c-0.39398 0.27742-0.79185 0.41643-1.1932 0.41643-0.63367 0-1.2033-0.18466-1.7093-0.55271-0.50611-0.36849-0.93275-0.876-1.2804-1.5229-0.34763-0.64644-0.61405-1.4063-0.79948-2.2793-0.18536-0.87282-0.27809-1.814-0.27809-2.8227 0-1.0341 0.09244-1.9946 0.27731-2.8806 0.18492-0.88574 0.45069-1.6583 0.7975-2.318 0.3469-0.65973 0.77306-1.1768 1.2788-1.5522 0.50576-0.37492 1.0754-0.5626 1.7091-0.5626 0.42477 0 0.83595 0.10702 1.2337 0.32078 0.39787 0.21405 0.75703 0.52833 1.0777 0.94299 0.32055 0.41458 0.58507 0.92571 0.79361 1.533 0.20862 0.60773 0.33627 1.3048 0.38258 2.092h-1.7379c-0.10825-0.77399-0.31688-1.3548-0.62585-1.7421-0.30918-0.38726-0.68377-0.58085-1.1239-0.58085-0.40953 0-0.75716 0.13251-1.0429 0.39767-0.28573 0.26516-0.51756 0.62102-0.69531 1.067-0.17779 0.44649-0.30703 0.95392-0.38794 1.5229-0.08131 0.5691-0.12163 1.1572-0.12163 1.765 0 0.58225 0.04057 1.148 0.12181 1.6975 0.08131 0.54938 0.21083 1.0441 0.38858 1.4838 0.17774 0.43955 0.40988 0.79172 0.69608 1.0569 0.28613 0.26516 0.63419 0.39774 1.0444 0.39774 0.60313 0 1.0688-0.25416 1.3976-0.76269 0.32884-0.50831 0.52028-1.2455 0.57449-2.2113h-1.8306v-2.2743h3.471v7.4747h-1.1549l-0.18497-1.5681c-0.32444 0.69696-0.68385 1.1844-1.0778 1.4619"/>
+ <path id="path24" d="m24.327 8.7799v11.255h4.0323v2.5658h-5.8471v-13.82h1.8149"/>
+ <path id="path26" d="m28.67 22.941c-2.3343 2.3722-5.8305 3.7224-9.7313 3.7224-7.028 0-12.725-4.8897-12.725-10.922 0-6.0319 5.6973-10.922 12.725-10.922 3.9152 0 7.4255 1.3753 9.7599 3.7623-2.548-3.9995-7.1164-6.8098-12.34-6.8111-8-0.0021579-14.486 6.251-14.487 13.967-7.113e-4 7.7155 6.484 13.972 14.484 13.974 5.2027 0.0015 9.7594-2.7998 12.314-6.7714"/>
+ </g>
+ <path id="path14-7" d="m7.072 17.859c0.10111 0.63058 0.25985 1.198 0.47648 1.7022 0.21643 0.50446 0.49426 0.90971 0.83362 1.2162 0.33907 0.30631 0.74316 0.45905 1.2125 0.45905 0.46911 0 0.87321-0.15274 1.2126-0.45905 0.33907-0.30652 0.6169-0.71176 0.83358-1.2162 0.21639-0.50426 0.37512-1.0717 0.47624-1.7022 0.10105-0.63064 0.15161-1.2702 0.15161-1.9186s-0.05057-1.2875-0.15161-1.9182c-0.10111-0.63064-0.25984-1.198-0.47624-1.7027-0.21667-0.50412-0.4945-0.90937-0.83358-1.2158-0.33935-0.30638-0.74345-0.4596-1.2126-0.4596-0.46936 0-0.87345 0.15322-1.2125 0.4596-0.33935 0.30645-0.61719 0.7117-0.83362 1.2158-0.21663 0.50467-0.37537 1.072-0.47648 1.7027-0.10107 0.63065-0.15137 1.2697-0.15137 1.9182s0.050295 1.288 0.15137 1.9186zm-0.94306-4.4517c0.15495-0.80974 0.38759-1.5269 0.69746-2.1507 0.30986-0.62383 0.69568-1.1221 1.1569-1.4938 0.46127-0.37181 0.99842-0.55813 1.6113-0.55813 0.6126 0 1.1498 0.18632 1.6111 0.55813 0.46123 0.37167 0.84705 0.86994 1.1569 1.4938 0.30986 0.62383 0.54251 1.341 0.69742 2.1507 0.15495 0.80974 0.23264 1.6526 0.23264 2.5285 0 0.87586-0.0777 1.7191-0.23264 2.5287-0.15491 0.80974-0.38756 1.5269-0.69742 2.1508-0.30986 0.62383-0.69568 1.119-1.1569 1.4848-0.46127 0.36548-0.99846 0.54877-1.6111 0.54877-0.61288 0-1.15-0.18329-1.6113-0.54877-0.46123-0.36582-0.84705-0.861-1.1569-1.4848-0.30986-0.6239-0.54251-1.3411-0.69746-2.1508-0.15491-0.8096-0.23241-1.6528-0.23241-2.5287 0-0.87593 0.0775-1.7188 0.23241-2.5285" fill="#4fa6db"/>
+</svg>
diff --git a/silx/resources/gui/icons/rm.png b/silx/resources/gui/icons/rm.png
new file mode 100644
index 0000000..ecff08b
--- /dev/null
+++ b/silx/resources/gui/icons/rm.png
Binary files 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 @@
+<?xml version="1.0" encoding="UTF-8"?>
+<svg id="svg2" enable-background="new 0 0 32 32" version="1.1" viewBox="0 0 32 32" xml:space="preserve" xmlns="http://www.w3.org/2000/svg" xmlns:cc="http://creativecommons.org/ns#" xmlns:dc="http://purl.org/dc/elements/1.1/" xmlns:rdf="http://www.w3.org/1999/02/22-rdf-syntax-ns#"><metadata id="metadata24"><rdf:RDF><cc:Work rdf:about=""><dc:format>image/svg+xml</dc:format><dc:type rdf:resource="http://purl.org/dc/dcmitype/StillImage"/><dc:title/></cc:Work></rdf:RDF></metadata><rect id="rect3059" x="3.993" y="4.378" width="23.404" height="22.837" color="#000000" enable-background="accumulate" fill="#fff" stroke="#000" stroke-miterlimit="10" stroke-width="1.8965"/><path id="path3877" d="m7.6413 16.32h15.571" color="#000000" enable-background="accumulate" fill="#fff" stroke="#000" stroke-miterlimit="10" stroke-width="3.9114"/></svg>
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<num_bins) {
- x = CLIP_MAX(h0, num_bins);
- FLOORCEIL_x(x);
- vals = ADJACENT_PIXELS_VALS(d_sino, num_bins, ym, xm, xp);
- res0 += linear_interpolation(vals, x, xm, xp);
- }
- if(h1>=0 && h1<num_bins) {
- x = CLIP_MAX(h1, num_bins);
- FLOORCEIL_x(x);
- vals = ADJACENT_PIXELS_VALS(d_sino, num_bins, ym, xm, xp);
- res1 += linear_interpolation(vals, x, xm, xp);
- }
- if(h2>=0 && h2<num_bins) {
- x = CLIP_MAX(h2, num_bins);
- FLOORCEIL_x(x);
- vals = ADJACENT_PIXELS_VALS(d_sino, num_bins, ym, xm, xp);
- res2 += linear_interpolation(vals, x, xm, xp);
- }
- if(h3>=0 && h3<num_bins) {
- x = CLIP_MAX(h3, num_bins);
- FLOORCEIL_x(x);
- vals = ADJACENT_PIXELS_VALS(d_sino, num_bins, ym, xm, xp);
- res3 += linear_interpolation(vals, x, xm, xp);
- }
- }
- d_SLICE[ 32*get_num_groups(0)*(bidy*32+tidy*2+0) + bidx*32 + tidx*2 + 0] = res0;
- d_SLICE[ 32*get_num_groups(0)*(bidy*32+tidy*2+1) + bidx*32 + tidx*2 + 0] = res1;
- 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;
-}
-
-
-
-
-
-
-
-/*******************************************************************************/
-/************************** OLD STUFF, for tinkering **************************/
-/*******************************************************************************/
-
-
-
-
-/// arr(xm, ym), arr(xm, yp), arr(xp, yp), arr(xp, ym)
-//~ #define ADJACENT_PIXELS_VALS2(arr, Nx, xm, xp, ym, yp) ((float4) (arr[ym*Nx + xm], arr[yp*Nx + xm], arr[yp*Nx + xp], arr[ym*Nx + xp]))
-
-
-/** xm, xp, ym, yp **/
-//~ #define ADJACENT_PIXELS_COORDS(x, y) ((int4)((int) floor(x), (int) ceil(x), (int) floor(y), (int) ceil(y)))
-
-/**
- (xm, ym) (xp, ym)
- (x, y)
- (xm, yp) (xp, yp)
-**/
-/// arr(xm, ym), arr(xm, yp), arr(xp, yp), arr(xp, ym)
-//~ #define ADJACENT_PIXELS_VALS(arr, Nx, coords) ((float4) (arr[coords.s2*Nx + coords.s0], arr[coords.s3*Nx + coords.s0], arr[coords.s3*Nx + coords.s1], arr[coords.s2*Nx + coords.s1]))
-
-
-/** xm, xp **/
-//~ #define ADJACENT_PIXELS_COORDS2(x) ((int2)((int) floor(x), (int) ceil(x)))
-
-
-
-/*
-float bilinear_interpolation(
- float x, // x position in the image
- float y, // y position in the image
- int Nx, // image width
- int Ny, // image height
- int4 adj_coords,
- float4 adj_vals
-) {
- float val;
- float tol = 0.001f; // CHECKME
- val = y - adj_coords.s2;
- if ((x - adj_coords.s0) < tol && (y - adj_coords.s2) < tol) val = adj_vals.s0;
- else if ((adj_coords.s1 - x) < tol && (adj_coords.s3 - y) < tol) val = adj_vals.s2;
- else {
- // Mirror - TODO: clamp ?
- if (adj_coords.s0 < 0) adj_coords.s0 = 0;
- if (adj_coords.s1 >= 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<num_proj; proj++) {
- if(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<num_bins) {
+ x = CLIP_MAX(h0, num_bins);
+ FLOORCEIL_x(x);
+ vals = ADJACENT_PIXELS_VALS(d_sino, num_bins, ym, xm, xp);
+ res0 += linear_interpolation(vals, x, xm, xp);
}
- pcos = sh_cos[256-read + proj] ;
- psin = sh_sin[256-read + proj] ;
-
- acorr05 = sh_axis[256 - read + proj] ;
-
- h0 = (acorr05 + bx00*pcos - by00*psin);
- 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, val;
- float tol = 0.001f; // CHECKME
- float y = proj + 0.5f;
- int ym = (int) floor(y);
- int yp = (int) ceil(y);
- int xm, xp;
-
- //
- int i0, i1, j0, j1;
- float d0, d1, x0, x1, y0, y1;
- d0 = fmin(fmax(proj+0*0.5f, 0.0f), (num_proj - 1.0f));
- x0 = floor(d0);
- x1 = ceil(d0);
- i0 = (int) x0;
- i1 = (int) x1;
-
- if(h0>=0 && h0<num_bins) {
- d1 = fmin(fmax(h0+0*0.5f, 0.0f), (num_bins - 1.0f));
- y0 = floor(d1);
- y1 = ceil(d1);
- j0 = (int) y0;
- j1 = (int) y1;
-
- if ((i0 == i1) && (j0 == j1))
- val = d_sino[i0*num_bins + j0]; //self.data[i0, j0]
- else if (i0 == i1)
- val = (d_sino[i0*num_bins + j0] * (y1 - d1)) + (d_sino[i0*num_bins + j1] * (d1 - y0)); // self.data[i0, j0], self.data[i0, j1]
- else if (j0 == j1)
- val = (d_sino[i0*num_bins + j0] * (x1 - d0)) + (d_sino[i1*num_bins + j0] * (d0 - x0)); // i0, j0 ; i1, j0
- else
- val = (d_sino[i0*num_bins + j0] * (x1 - d0) * (y1 - d1)) // i0, j0
- + (d_sino[i1*num_bins + j0] * (d0 - x0) * (y1 - d1)) // i1, j0
- + (d_sino[i0*num_bins + j1] * (x1 - d0) * (d1 - y0)) // i0, j1
- + (d_sino[i1*num_bins + j1] * (d0 - x0) * (d1 - y0)); // i1, j1
-
- res0 += val;
- }
if(h1>=0 && h1<num_bins) {
- //~ int4 coords = ADJACENT_PIXELS_COORDS(h1 +0.5f, proj +0.5f);
- //~ res1 += bilinear_interpolation(h1 +0.5f, proj +0.5f, num_bins, num_proj, coords, ADJACENT_PIXELS_VALS(d_sino, num_bins, coords)); //tex2D(texProjes,h1 +0.5f,proj +0.5f);
- d1 = fmin(fmax(h1+0*0.5f, 0.0f), (num_bins - 1.0f));
- y0 = floor(d1);
- y1 = ceil(d1);
- j0 = (int) y0;
- j1 = (int) y1;
-
- if ((i0 == i1) && (j0 == j1))
- val = d_sino[i0*num_bins + j0]; //self.data[i0, j0]
- else if (i0 == i1)
- val = (d_sino[i0*num_bins + j0] * (y1 - d1)) + (d_sino[i0*num_bins + j1] * (d1 - y0)); // self.data[i0, j0], self.data[i0, j1]
- else if (j0 == j1)
- val = (d_sino[i0*num_bins + j0] * (x1 - d0)) + (d_sino[i1*num_bins + j0] * (d0 - x0)); // i0, j0 ; i1, j0
- else
- val = (d_sino[i0*num_bins + j0] * (x1 - d0) * (y1 - d1)) // i0, j0
- + (d_sino[i1*num_bins + j0] * (d0 - x0) * (y1 - d1)) // i1, j0
- + (d_sino[i0*num_bins + j1] * (x1 - d0) * (d1 - y0)) // i0, j1
- + (d_sino[i1*num_bins + j1] * (d0 - x0) * (d1 - y0)); // i1, j1
-
-
- res1 += val;
+ x = CLIP_MAX(h1, num_bins);
+ FLOORCEIL_x(x);
+ vals = ADJACENT_PIXELS_VALS(d_sino, num_bins, ym, xm, xp);
+ res1 += linear_interpolation(vals, x, xm, xp);
}
if(h2>=0 && h2<num_bins) {
- //~ int4 coords = ADJACENT_PIXELS_COORDS(h2 +0.5f, proj +0.5f);
- //~ res2 += 0; //bilinear_interpolation(h2 +0.5f, proj +0.5f, num_bins, num_proj, coords, ADJACENT_PIXELS_VALS(d_sino, num_bins, coords)); //tex2D(texProjes,h2 +0.5f,proj +0.5f);
- d1 = fmin(fmax(h2+0*0.5f, 0.0f), (num_bins - 1.0f));
- y0 = floor(d1);
- y1 = ceil(d1);
- j0 = (int) y0;
- j1 = (int) y1;
-
- if ((i0 == i1) && (j0 == j1))
- val = d_sino[i0*num_bins + j0]; //self.data[i0, j0]
- else if (i0 == i1)
- val = (d_sino[i0*num_bins + j0] * (y1 - d1)) + (d_sino[i0*num_bins + j1] * (d1 - y0)); // self.data[i0, j0], self.data[i0, j1]
- else if (j0 == j1)
- val = (d_sino[i0*num_bins + j0] * (x1 - d0)) + (d_sino[i1*num_bins + j0] * (d0 - x0)); // i0, j0 ; i1, j0
- else
- val = (d_sino[i0*num_bins + j0] * (x1 - d0) * (y1 - d1)) // i0, j0
- + (d_sino[i1*num_bins + j0] * (d0 - x0) * (y1 - d1)) // i1, j0
- + (d_sino[i0*num_bins + j1] * (x1 - d0) * (d1 - y0)) // i0, j1
- + (d_sino[i1*num_bins + j1] * (d0 - x0) * (d1 - y0)); // i1, j1
-
- res2+= val;
+ x = CLIP_MAX(h2, num_bins);
+ FLOORCEIL_x(x);
+ vals = ADJACENT_PIXELS_VALS(d_sino, num_bins, ym, xm, xp);
+ res2 += linear_interpolation(vals, x, xm, xp);
}
if(h3>=0 && h3<num_bins) {
- //~ int4 coords = ADJACENT_PIXELS_COORDS(h3 +0.5f, proj +0.5f);
- //~ res3 += 0; //bilinear_interpolation(h3 +0.5f, proj +0.5f, num_bins, num_proj, coords, ADJACENT_PIXELS_VALS(d_sino, num_bins, coords)); //tex2D(texProjes,h3 +0.5f,proj +0.5f);
- d1 = fmin(fmax(h3+0*0.5f, 0.0f), (num_bins - 1.0f));
- y0 = floor(d1);
- y1 = ceil(d1);
- j0 = (int) y0;
- j1 = (int) y1;
-
- if ((i0 == i1) && (j0 == j1))
- val = d_sino[i0*num_bins + j0]; //self.data[i0, j0]
- else if (i0 == i1)
- val = (d_sino[i0*num_bins + j0] * (y1 - d1)) + (d_sino[i0*num_bins + j1] * (d1 - y0)); // self.data[i0, j0], self.data[i0, j1]
- else if (j0 == j1)
- val = (d_sino[i0*num_bins + j0] * (x1 - d0)) + (d_sino[i1*num_bins + j0] * (d0 - x0)); // i0, j0 ; i1, j0
- else
- val = (d_sino[i0*num_bins + j0] * (x1 - d0) * (y1 - d1)) // i0, j0
- + (d_sino[i1*num_bins + j0] * (d0 - x0) * (y1 - d1)) // i1, j0
- + (d_sino[i0*num_bins + j1] * (x1 - d0) * (d1 - y0)) // i0, j1
- + (d_sino[i1*num_bins + j1] * (d0 - x0) * (d1 - y0)); // i1, j1
-
- res3 += val;
+ x = CLIP_MAX(h3, num_bins);
+ FLOORCEIL_x(x);
+ vals = ADJACENT_PIXELS_VALS(d_sino, num_bins, ym, xm, xp);
+ res3 += linear_interpolation(vals, x, xm, xp);
}
}
d_SLICE[ 32*get_num_groups(0)*(bidy*32+tidy*2+0) + bidx*32 + tidx*2 + 0] = res0;
@@ -478,8 +229,4 @@ __kernel void backproj_cpu_kernel_good(
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;
}
-*/
-
-
-
diff --git a/silx/resources/opencl/proj.cl b/silx/resources/opencl/proj.cl
index afc58ff..2a6d870 100644
--- a/silx/resources/opencl/proj.cl
+++ b/silx/resources/opencl/proj.cl
@@ -28,7 +28,7 @@
/************************ GPU VERSION (with textures) **************************/
/*******************************************************************************/
-
+#ifndef DONT_USE_TEXTURES
kernel void forward_kernel(
global float *d_Sino,
read_only image2d_t d_slice,
@@ -163,7 +163,7 @@ kernel void forward_kernel(
d_Sino[dimrecx*(bidy*16 + tidy) + (bidx*16 + tidx)] = res;
}
}
-
+#endif
/*******************************************************************************/
generated by cgit v1.2.3 (git 2.25.1) at 2024-04-28 17:47:08 +0000