CustusX  2023.01.05-dev+develop.0da12
An IGT application
Classes | Macros | Typedefs | Functions
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 137 of file kernels.cl.h.

#define ANISOTROPIC_GAUSSIAN_SIZE   3

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

#define ANISOTROPIC_METHOD   ANISOTROPIC_METHOD_BILINEAR_ON_PLANE

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

#define ANISOTROPIC_METHOD_BILINEAR_ON_PLANE   1

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

#define ANISOTROPIC_METHOD_GAUSSIAN_ON_PLANE   2

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

#define ANISOTROPIC_METHOD_INCLUDE_ALL   0

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

#define ANISOTROPIC_SIZE   3

Definition at line 53 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 169 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)))
BRIGHTNESS_FACTOR
#define BRIGHTNESS_FACTOR
Definition: kernels.cl.h:144
WEIGHT_TERNARY
#define WEIGHT_TERNARY(val, mean, factor)
Definition: kernels.cl.h:134
NEWNESS_FACTOR
#define NEWNESS_FACTOR
Definition: kernels.cl.h:148
WEIGHT_GAUSS
#define WEIGHT_GAUSS(x, sigma)
Definition: kernels.cl.h:180

Definition at line 157 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 151 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 154 of file kernels.cl.h.

#define ANISOTROPIC_WEIGHT_METHOD   ANISOTROPIC_WEIGHT_METHOD_BOTH

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

#define ANISOTROPIC_WEIGHT_METHOD_BOTH   3

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

#define ANISOTROPIC_WEIGHT_METHOD_BRIGHTNESS   1

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

#define ANISOTROPIC_WEIGHT_METHOD_DISTANCE   0

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

#define ANISOTROPIC_WEIGHT_METHOD_LATENESS   2

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

#define BOUNDS_CHECK (   x,
  min,
  max 
)

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

#define BRIGHTNESS_FACTOR   5.0f

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

#define CHECK_PLANE_INDICES

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

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

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

#define CREATE_OUTPUT_VOLUME_TYPE (   name,
  in_size,
  in_spacing,
  in_volume 
)
Value:
output_volume_type name; \
name.size = in_size; \
name.spacing = in_spacing; \
name.volume = in_volume;
_output_volume_type
Definition: kernels.cl.h:87
_output_volume_type::size
int3 size
Definition: kernels.cl.h:89

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

#define CUBE_SIZE   4

Original author Tord �ygard

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

#define DW_WEIGHT (   x)    WEIGHT_INV(x)

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

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

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

#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)

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

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

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

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

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

#define METHOD_ANISOTROPIC   3

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

#define METHOD_DW   2

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

#define METHOD_VNN   0

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

#define METHOD_VNN2   1

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

#define N_BLOCKS   10

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

#define NEWNESS_FACTOR   5.0f

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

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

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

#define PLANE_CLOSEST   1

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

#define PLANE_HEURISTIC   0

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

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

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

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

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

#define VNN2_WEIGHT (   x)    WEIGHT_INV(x)

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

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

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

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

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

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

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

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

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

#define WEIGHT_GAUSS_SIGMA   (0.05f)

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

#define WEIGHT_GAUSS_SQRT_2PI   2.506628275f

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

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

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

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

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

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

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

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

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

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

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

Typedef Documentation

typedef struct _close_plane close_plane_t
typedef struct _output_volume_type output_volume_type

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 
)