NorMIT-nav  16.5
An IGT application
 All Classes Namespaces Files Functions Variables Typedefs Enumerations Enumerator Friends Macros Groups Pages
kernels.cl.h File Reference

Go to the source code of this file.

Classes

struct  _close_plane
 
struct  _output_volume_type
 

Macros

#define CUBE_SIZE   4
 
#define N_BLOCKS   10
 
#define METHOD_VNN   0
 
#define METHOD_VNN2   1
 
#define METHOD_DW   2
 
#define METHOD_ANISOTROPIC   3
 
#define PLANE_HEURISTIC   0
 
#define PLANE_CLOSEST   1
 
#define ANISOTROPIC_SIZE   3
 
#define ANISOTROPIC_GAUSSIAN_SIZE   3
 
#define ANISOTROPIC_METHOD_INCLUDE_ALL   0
 
#define ANISOTROPIC_METHOD_BILINEAR_ON_PLANE   1
 
#define ANISOTROPIC_METHOD_GAUSSIAN_ON_PLANE   2
 
#define ANISOTROPIC_WEIGHT_METHOD_DISTANCE   0
 
#define ANISOTROPIC_WEIGHT_METHOD_BRIGHTNESS   1
 
#define ANISOTROPIC_WEIGHT_METHOD_LATENESS   2
 
#define ANISOTROPIC_WEIGHT_METHOD_BOTH   3
 
#define ANISOTROPIC_METHOD   ANISOTROPIC_METHOD_BILINEAR_ON_PLANE
 
#define CHECK_PLANE_INDICES
 
#define BOUNDS_CHECK(x, min, max)
 
#define EUCLID_DIST(a, b, c)   sqrt((a)*(a) + (b)*(b) + (c)*(c))
 
#define PROJECTONTOPLANE(voxel, matrix, dist)   (voxel - dist*(matrix.s26AE))
 
#define PROJECTONTOPLANEEQ(voxel, eq, dist)   (voxel - dist*(eq))
 
#define ISINSIDE(x, size)   ((x) >= 0 && (x) < (size))
 
#define ISNOTMASKED(x, y, mask, xsize)   ((mask)[(x) + (y)*(xsize)] > 0)
 
#define VOXEL(v, x, y, z)   v[x + y*volume_xsize + z*volume_ysize*volume_xsize]
 
#define WEIGHT_INV(x)   (1.0f/fabs(x))
 
#define WEIGHT_INV2(x)   (1.0f/fabs(x*x))
 
#define WEIGHT_INV4(x)   (1.0f/fabs(x*x*x*x))
 
#define WEIGHT_SUB(x)   (1.0f - fabs(x))
 
#define WEIGHT_TERNARY(val, mean, factor)   ((val) >= (mean) ? (factor) : 0.0f)
 
#define ANISOTROPIC_GAUSS_WEIGHT(px, var, mean, mean_id, sigma)   WEIGHT_GAUSS(px.dist, sigma)
 
#define ANISOTROPIC_WEIGHT_METHOD   ANISOTROPIC_WEIGHT_METHOD_BOTH
 
#define BRIGHTNESS_FACTOR   5.0f
 
#define NEWNESS_FACTOR   5.0f
 
#define ANISOTROPIC_WEIGHT_BRIGHTNESS(px, var, mean, mean_id, sigma)   ((WEIGHT_GAUSS(px.dist, sigma)) + (WEIGHT_TERNARY(px.intensity, mean, BRIGHTNESS_FACTOR)))
 
#define ANISOTROPIC_WEIGHT_LATENESS(px, var, mean, mean_id, sigma)   ((WEIGHT_GAUSS(px.dist, sigma)) + (WEIGHT_TERNARY(px.plane_id, mean_id, NEWNESS_FACTOR)))
 
#define ANISOTROPIC_WEIGHT_BOTH(px, var, mean, mean_id, sigma)
 
#define ANISOTROPIC_WEIGHT(px, var, mean, mean_id, sigma)   ANISOTROPIC_WEIGHT_BOTH(px, var, mean, mean_id, sigma)
 
#define WEIGHT_GAUSS_SIGMA   (0.05f)
 
#define WEIGHT_GAUSS_SQRT_2PI   2.506628275f
 
#define WEIGHT_GAUSS_NONEXP_PART(sigma)   (1.0f/((sigma)*WEIGHT_GAUSS_SQRT_2PI))
 
#define WEIGHT_GAUSS_EXP_PART(dist, sigma)   exp(-((dist)*(dist))/(2*(sigma)*(sigma)))
 
#define WEIGHT_GAUSS(x, sigma)   (WEIGHT_GAUSS_NONEXP_PART(sigma)*WEIGHT_GAUSS_EXP_PART(x, sigma))
 
#define DW_WEIGHT(x)   WEIGHT_INV(x)
 
#define VNN2_WEIGHT(x)   WEIGHT_INV(x)
 
#define CLOSE_PLANE_IDX(p, i)   p[get_local_id(0)*(MAX_PLANES+1)+(i)]
 
#define CREATE_OUTPUT_VOLUME_TYPE(name, in_size, in_spacing, in_volume)
 
#define FIND_CLOSE_PLANES(a, b, c, d, e, f, g, h, i, j)   findClosestPlanes_multistart(a, b, c, d, e, f, g, 1, h, i, j)
 
#define PERFORM_INTERPOLATION(a, b, c, d, e, f, g, h, i)   performInterpolation_vnn(a, b, c, d, e, f, g, h, i)
 

Typedefs

typedef struct _close_plane close_plane_t
 
typedef struct _output_volume_type output_volume_type
 

Functions

int isValidPixel (int2 point, const __global unsigned char *mask, int2 in_size)
 
int findHighestIdx (__local close_plane_t *planes, int n)
 
int2 findClosestPlanes_heuristic (__local close_plane_t *close_planes, __local float4 *const plane_eqs, __global float16 *const plane_matrices, const float4 voxel, const float radius, int guess, bool doTermDistance, __global const unsigned char *mask, int2 in_size, float2 in_spacing)
 
int2 findClosestPlanes_multistart (__local close_plane_t *close_planes, __local float4 *const plane_eqs, __global float16 *const plane_matrices, const float4 voxel, const float radius, int *multistart_guesses, int n_multistart_guesses, bool doTermDistance, __global const unsigned char *mask, int2 in_size, float2 in_spacing)
 
__global const unsigned char * getImageData (int plane_id, __global const unsigned char *bscans_blocks[], int2 in_size)
 
float4 transform (float16 matrix, float4 voxel)
 
float4 transform_inv (float16 matrix, float4 voxel)
 
float2 transform_inv_xy (float16 matrix, float4 voxel)
 
int2 round_int (float2 value)
 
int2 toImgCoord_int (float4 voxel, float16 plane_matrix, float2 in_spacing)
 
float2 toImgCoord_float (float4 voxel, float16 plane_matrix, float2 in_spacing)
 
float bilinearInterpolation (float x, float y, const __global unsigned char *image, int in_xsize)
 
unsigned char performInterpolation_vnn (__local close_plane_t *close_planes, int n_close_planes, __global const float16 *plane_matrices, __local const float4 *plane_eqs, __global const unsigned char *bscans_blocks[], int2 in_size, float2 in_spacing, __global const unsigned char *mask, float4 voxel)
 
unsigned char performInterpolation_vnn2 (__local close_plane_t *close_planes, int n_close_planes, __global const float16 *plane_matrices, __local const float4 *plane_eqs, __global const unsigned char *bscans_blocks[], int2 in_size, float2 in_spacing, __global const unsigned char *mask, float4 voxel)
 
unsigned char performInterpolation_dw (__local close_plane_t *close_planes, int n_close_planes, __global const float16 *plane_matrices, __local const float4 *plane_eqs, __global const unsigned char *bscans_blocks[], int2 in_size, float2 in_spacing, __global const unsigned char *mask, float4 voxel)
 
unsigned char performInterpolation_anisotropic (__local close_plane_t *close_planes, int n_close_planes, __global const float16 *plane_matrices, __local const float4 *plane_eqs, __global const unsigned char *bscans_blocks[], int2 in_size, float2 in_spacing, __global const unsigned char *mask, float4 voxel)
 
unsigned char anisotropicFilter (__local const close_plane_t *pixels, int n_planes)
 
void prepare_plane_eqs (__global float16 *plane_matrices, __local float4 *plane_eqs)
 
int findLocalMinimas (int *guesses, __local const float4 *plane_eqs, float radius, float4 voxel, float3 out_spacing, __global const float16 *plane_matrices, __global const unsigned char *mask)
 
__kernel void voxel_methods (int volume_xsize, int volume_ysize, int volume_zsize, float volume_xspacing, float volume_yspacing, float volume_zspacing, int in_xsize, int in_ysize, float in_xspacing, float in_yspacing, __global unsigned char *in_bscans_b0, __global unsigned char *in_bscans_b1, __global unsigned char *in_bscans_b2, __global unsigned char *in_bscans_b3, __global unsigned char *in_bscans_b4, __global unsigned char *in_bscans_b5, __global unsigned char *in_bscans_b6, __global unsigned char *in_bscans_b7, __global unsigned char *in_bscans_b8, __global unsigned char *in_bscans_b9, __global unsigned char *out_volume, __global float16 *plane_matrices, __global unsigned char *mask, __local float4 *plane_eqs, __local close_plane_t *planes, float radius)
 

Macro Definition Documentation

#define ANISOTROPIC_GAUSS_WEIGHT (   px,
  var,
  mean,
  mean_id,
  sigma 
)    WEIGHT_GAUSS(px.dist, sigma)

Definition at line 158 of file kernels.cl.h.

#define ANISOTROPIC_GAUSSIAN_SIZE   3

Definition at line 75 of file kernels.cl.h.

#define ANISOTROPIC_METHOD   ANISOTROPIC_METHOD_BILINEAR_ON_PLANE

Definition at line 86 of file kernels.cl.h.

#define ANISOTROPIC_METHOD_BILINEAR_ON_PLANE   1

Definition at line 78 of file kernels.cl.h.

#define ANISOTROPIC_METHOD_GAUSSIAN_ON_PLANE   2

Definition at line 79 of file kernels.cl.h.

#define ANISOTROPIC_METHOD_INCLUDE_ALL   0

Definition at line 77 of file kernels.cl.h.

#define ANISOTROPIC_SIZE   3

Definition at line 74 of file kernels.cl.h.

#define ANISOTROPIC_WEIGHT (   px,
  var,
  mean,
  mean_id,
  sigma 
)    ANISOTROPIC_WEIGHT_BOTH(px, var, mean, mean_id, sigma)

Definition at line 190 of file kernels.cl.h.

#define ANISOTROPIC_WEIGHT_BOTH (   px,
  var,
  mean,
  mean_id,
  sigma 
)
Value:
((WEIGHT_GAUSS(px.dist, sigma)) \
+ (WEIGHT_TERNARY(px.plane_id, mean_id, NEWNESS_FACTOR)) \
+ (WEIGHT_TERNARY(px.intensity, mean, BRIGHTNESS_FACTOR)))
#define BRIGHTNESS_FACTOR
Definition: kernels.cl.h:165
#define WEIGHT_TERNARY(val, mean, factor)
Definition: kernels.cl.h:155
#define NEWNESS_FACTOR
Definition: kernels.cl.h:169
#define WEIGHT_GAUSS(x, sigma)
Definition: kernels.cl.h:201

Definition at line 178 of file kernels.cl.h.

#define ANISOTROPIC_WEIGHT_BRIGHTNESS (   px,
  var,
  mean,
  mean_id,
  sigma 
)    ((WEIGHT_GAUSS(px.dist, sigma)) + (WEIGHT_TERNARY(px.intensity, mean, BRIGHTNESS_FACTOR)))

Definition at line 172 of file kernels.cl.h.

#define ANISOTROPIC_WEIGHT_LATENESS (   px,
  var,
  mean,
  mean_id,
  sigma 
)    ((WEIGHT_GAUSS(px.dist, sigma)) + (WEIGHT_TERNARY(px.plane_id, mean_id, NEWNESS_FACTOR)))

Definition at line 175 of file kernels.cl.h.

#define ANISOTROPIC_WEIGHT_METHOD   ANISOTROPIC_WEIGHT_METHOD_BOTH

Definition at line 161 of file kernels.cl.h.

#define ANISOTROPIC_WEIGHT_METHOD_BOTH   3

Definition at line 84 of file kernels.cl.h.

#define ANISOTROPIC_WEIGHT_METHOD_BRIGHTNESS   1

Definition at line 82 of file kernels.cl.h.

#define ANISOTROPIC_WEIGHT_METHOD_DISTANCE   0

Definition at line 81 of file kernels.cl.h.

#define ANISOTROPIC_WEIGHT_METHOD_LATENESS   2

Definition at line 83 of file kernels.cl.h.

#define BOUNDS_CHECK (   x,
  min,
  max 
)

Definition at line 133 of file kernels.cl.h.

#define BRIGHTNESS_FACTOR   5.0f

Definition at line 165 of file kernels.cl.h.

#define CHECK_PLANE_INDICES

Definition at line 124 of file kernels.cl.h.

#define CLOSE_PLANE_IDX (   p,
 
)    p[get_local_id(0)*(MAX_PLANES+1)+(i)]

Definition at line 206 of file kernels.cl.h.

#define CREATE_OUTPUT_VOLUME_TYPE (   name,
  in_size,
  in_spacing,
  in_volume 
)
Value:
name.size = in_size; \
name.spacing = in_spacing; \
name.volume = in_volume;

Definition at line 208 of file kernels.cl.h.

#define CUBE_SIZE   4

Original author Tord �ygard

Definition at line 59 of file kernels.cl.h.

#define DW_WEIGHT (   x)    WEIGHT_INV(x)

Definition at line 203 of file kernels.cl.h.

#define EUCLID_DIST (   a,
  b,
 
)    sqrt((a)*(a) + (b)*(b) + (c)*(c))

Definition at line 138 of file kernels.cl.h.

#define FIND_CLOSE_PLANES (   a,
  b,
  c,
  d,
  e,
  f,
  g,
  h,
  i,
 
)    findClosestPlanes_multistart(a, b, c, d, e, f, g, 1, h, i, j)

Definition at line 258 of file kernels.cl.h.

#define ISINSIDE (   x,
  size 
)    ((x) >= 0 && (x) < (size))

Definition at line 144 of file kernels.cl.h.

#define ISNOTMASKED (   x,
  y,
  mask,
  xsize 
)    ((mask)[(x) + (y)*(xsize)] > 0)

Definition at line 145 of file kernels.cl.h.

#define METHOD_ANISOTROPIC   3

Definition at line 67 of file kernels.cl.h.

#define METHOD_DW   2

Definition at line 66 of file kernels.cl.h.

#define METHOD_VNN   0

Definition at line 64 of file kernels.cl.h.

#define METHOD_VNN2   1

Definition at line 65 of file kernels.cl.h.

#define N_BLOCKS   10

Definition at line 61 of file kernels.cl.h.

#define NEWNESS_FACTOR   5.0f

Definition at line 169 of file kernels.cl.h.

#define PERFORM_INTERPOLATION (   a,
  b,
  c,
  d,
  e,
  f,
  g,
  h,
 
)    performInterpolation_vnn(a, b, c, d, e, f, g, h, i)

Definition at line 290 of file kernels.cl.h.

#define PLANE_CLOSEST   1

Definition at line 71 of file kernels.cl.h.

#define PLANE_HEURISTIC   0

Definition at line 70 of file kernels.cl.h.

#define PROJECTONTOPLANE (   voxel,
  matrix,
  dist 
)    (voxel - dist*(matrix.s26AE))

Definition at line 140 of file kernels.cl.h.

#define PROJECTONTOPLANEEQ (   voxel,
  eq,
  dist 
)    (voxel - dist*(eq))

Definition at line 142 of file kernels.cl.h.

#define VNN2_WEIGHT (   x)    WEIGHT_INV(x)

Definition at line 204 of file kernels.cl.h.

#define VOXEL (   v,
  x,
  y,
 
)    v[x + y*volume_xsize + z*volume_ysize*volume_xsize]

Definition at line 148 of file kernels.cl.h.

#define WEIGHT_GAUSS (   x,
  sigma 
)    (WEIGHT_GAUSS_NONEXP_PART(sigma)*WEIGHT_GAUSS_EXP_PART(x, sigma))

Definition at line 201 of file kernels.cl.h.

#define WEIGHT_GAUSS_EXP_PART (   dist,
  sigma 
)    exp(-((dist)*(dist))/(2*(sigma)*(sigma)))

Definition at line 199 of file kernels.cl.h.

#define WEIGHT_GAUSS_NONEXP_PART (   sigma)    (1.0f/((sigma)*WEIGHT_GAUSS_SQRT_2PI))

Definition at line 198 of file kernels.cl.h.

#define WEIGHT_GAUSS_SIGMA   (0.05f)

Definition at line 194 of file kernels.cl.h.

#define WEIGHT_GAUSS_SQRT_2PI   2.506628275f

Definition at line 196 of file kernels.cl.h.

#define WEIGHT_INV (   x)    (1.0f/fabs(x))

Definition at line 150 of file kernels.cl.h.

#define WEIGHT_INV2 (   x)    (1.0f/fabs(x*x))

Definition at line 151 of file kernels.cl.h.

#define WEIGHT_INV4 (   x)    (1.0f/fabs(x*x*x*x))

Definition at line 152 of file kernels.cl.h.

#define WEIGHT_SUB (   x)    (1.0f - fabs(x))

Definition at line 153 of file kernels.cl.h.

#define WEIGHT_TERNARY (   val,
  mean,
  factor 
)    ((val) >= (mean) ? (factor) : 0.0f)

Definition at line 155 of file kernels.cl.h.

Typedef Documentation

typedef struct _close_plane close_plane_t

Function Documentation

unsigned char anisotropicFilter ( __local const close_plane_t pixels,
int  n_planes 
)
float bilinearInterpolation ( float  x,
float  y,
const __global unsigned char *  image,
int  in_xsize 
)
int2 findClosestPlanes_heuristic ( __local close_plane_t close_planes,
__local float4 *const  plane_eqs,
__global float16 *const  plane_matrices,
const float4  voxel,
const float  radius,
int  guess,
bool  doTermDistance,
__global const unsigned char *  mask,
int2  in_size,
float2  in_spacing 
)
int2 findClosestPlanes_multistart ( __local close_plane_t close_planes,
__local float4 *const  plane_eqs,
__global float16 *const  plane_matrices,
const float4  voxel,
const float  radius,
int *  multistart_guesses,
int  n_multistart_guesses,
bool  doTermDistance,
__global const unsigned char *  mask,
int2  in_size,
float2  in_spacing 
)
int findHighestIdx ( __local close_plane_t planes,
int  n 
)
int findLocalMinimas ( int *  guesses,
__local const float4 *  plane_eqs,
float  radius,
float4  voxel,
float3  out_spacing,
__global const float16 *  plane_matrices,
__global const unsigned char *  mask 
)
__global const unsigned char* getImageData ( int  plane_id,
__global const unsigned char *  bscans_blocks[],
int2  in_size 
)
int isValidPixel ( int2  point,
const __global unsigned char *  mask,
int2  in_size 
)
unsigned char performInterpolation_anisotropic ( __local close_plane_t close_planes,
int  n_close_planes,
__global const float16 *  plane_matrices,
__local const float4 *  plane_eqs,
__global const unsigned char *  bscans_blocks[],
int2  in_size,
float2  in_spacing,
__global const unsigned char *  mask,
float4  voxel 
)
unsigned char performInterpolation_dw ( __local close_plane_t close_planes,
int  n_close_planes,
__global const float16 *  plane_matrices,
__local const float4 *  plane_eqs,
__global const unsigned char *  bscans_blocks[],
int2  in_size,
float2  in_spacing,
__global const unsigned char *  mask,
float4  voxel 
)
unsigned char performInterpolation_vnn ( __local close_plane_t close_planes,
int  n_close_planes,
__global const float16 *  plane_matrices,
__local const float4 *  plane_eqs,
__global const unsigned char *  bscans_blocks[],
int2  in_size,
float2  in_spacing,
__global const unsigned char *  mask,
float4  voxel 
)
unsigned char performInterpolation_vnn2 ( __local close_plane_t close_planes,
int  n_close_planes,
__global const float16 *  plane_matrices,
__local const float4 *  plane_eqs,
__global const unsigned char *  bscans_blocks[],
int2  in_size,
float2  in_spacing,
__global const unsigned char *  mask,
float4  voxel 
)
void prepare_plane_eqs ( __global float16 *  plane_matrices,
__local float4 *  plane_eqs 
)
int2 round_int ( float2  value)
float2 toImgCoord_float ( float4  voxel,
float16  plane_matrix,
float2  in_spacing 
)
int2 toImgCoord_int ( float4  voxel,
float16  plane_matrix,
float2  in_spacing 
)
float4 transform ( float16  matrix,
float4  voxel 
)
float4 transform_inv ( float16  matrix,
float4  voxel 
)
float2 transform_inv_xy ( float16  matrix,
float4  voxel 
)
__kernel void voxel_methods ( int  volume_xsize,
int  volume_ysize,
int  volume_zsize,
float  volume_xspacing,
float  volume_yspacing,
float  volume_zspacing,
int  in_xsize,
int  in_ysize,
float  in_xspacing,
float  in_yspacing,
__global unsigned char *  in_bscans_b0,
__global unsigned char *  in_bscans_b1,
__global unsigned char *  in_bscans_b2,
__global unsigned char *  in_bscans_b3,
__global unsigned char *  in_bscans_b4,
__global unsigned char *  in_bscans_b5,
__global unsigned char *  in_bscans_b6,
__global unsigned char *  in_bscans_b7,
__global unsigned char *  in_bscans_b8,
__global unsigned char *  in_bscans_b9,
__global unsigned char *  out_volume,
__global float16 *  plane_matrices,
__global unsigned char *  mask,
__local float4 *  plane_eqs,
__local close_plane_t planes,
float  radius 
)