1234567891011121314151617181920212223242526272829303132333435363738394041424344454647484950515253545556575859606162636465666768697071727374757677787980818283848586878889909192939495969798 |
- // all kernels must process volme voxelwise
- // please be careful with local and global workers
- // we need a 3D processing (not 2D)
- // using tests we show max volume grosse: 1024 x 1024 x 256 or 2048 x 2048 x 64
- // theory: 223 MB (CL_DEVICE_MAX_MEM_ALLOC_SIZE) x 5 cards / 4 Bytes (Float) = 1024 x 1024 x 278 voxels
- #include <lamino-filter-def.h>
- __kernel void lamino_bp_generic ( __global float *proj,
- __global float *volume,
- __global CLParameters *param)
- {
- const int vX = get_global_id(0);
- const int vY = get_global_id(1);
- const int vZ = get_global_id(2);
- const int vSX = get_global_size(0);
- const int vSY = get_global_size(1);
- const int vSZ = get_global_size(2);
- const long int idx = (vZ * vSY * vSX) + (vY * vSX) + vX;
- //const int idx = (vY * vSX) + vX;
- float newz = (float)vZ - param->vol_oz;
- float newz_matr02 = newz * param->mat_2 + param->proj_ox;
- float newz_matr12 = newz * param->mat_5 + param->proj_oy;
- /// prepare y info
- float newy = (float)vY - param->vol_oy;
- float newy_matr01 = newy * param->mat_1 + newz_matr02;
- float newy_matr11 = newy * param->mat_4 + newz_matr12;
- /// prepare x info
- float newx = (float)vX - param->vol_ox;
- float oldy = newx * param->mat_3 + newy_matr11;
- float yo = floor(oldy);
- float oldx = newx*param->mat_0 + newy_matr01;
- float xo = floor(oldx);
- // bilinear interpolation
- float yf_1 = oldy - yo;
- float yf_0 = 1.0f - yf_1;
- float xf_1 = oldx - xo;
- float xf_0 = 1.0f - xf_1;
- int base = (int)xo + (int)yo * param->proj_sx;
- float result;
- // TODO: check that out of proj plain
- result = proj[base ] * xf_0 * yf_0;
- result += proj[base + 1] * xf_1 * yf_0;
- result += proj[base + param->proj_sx ] * xf_0 * yf_1;
- result += proj[base + param->proj_sx + 1] * xf_1 * yf_1;
- volume[idx] += result;
- }
- __kernel void lamino_clean_vol(__global float *volume)
- {
- const int vX = get_global_id(0);
- const int vY = get_global_id(1);
- const int vZ = get_global_id(2);
- const int vSX = get_global_size(0);
- const int vSY = get_global_size(1);
- const int vSZ = get_global_size(2);
- const int idx = (vZ * vSY * vSX) + (vY * vSX) + vX;
- //const int idx = (vY * vSX) + vX;
- volume[idx] = 0;
- }
- __kernel void lamino_norm_vol(__global float *volume,
- const float factor)
- {
- const int vX = get_global_id(0);
- const int vY = get_global_id(1);
- const int vZ = get_global_id(2);
- const int vSX = get_global_size(0);
- const int vSY = get_global_size(1);
- const int vSZ = get_global_size(2);
- const int idx = (vZ * vSY * vSX) + (vY * vSX) + vX;
- //const int idx = (vY * vSX) + vX;
- float val = volume[idx] * factor;
- volume[idx] = val;
- }
|