Optimized Belief Propagation (CPU and GPU)
Functions
beliefprop_cuda Namespace Reference

Namespace to define global kernel functions for parallel belief propagation processing using CUDA. More...

Functions

template<RunData_t T, unsigned int DISP_VALS>
__global__ void InitializeBottomLevelData (beliefprop::BpLevelProperties current_bp_level, float *image_1_pixels_device, float *image_2_pixels_device, T *data_cost_stereo_checkerboard_0, T *data_cost_stereo_checkerboard_1, float lambda_bp, float data_k_bp, unsigned int bp_settings_disp_vals)
 Initialize the "data cost" for each possible disparity between the two full-sized input images ("bottom" of the image pyramid). The image data is stored in the image_1_pixels_device and image_2_pixels_device arrays. More...
 
template<RunData_t T, unsigned int DISP_VALS>
__global__ void InitializeCurrentLevelData (beliefprop::CheckerboardPart checkerboard_part, beliefprop::BpLevelProperties current_bp_level, beliefprop::BpLevelProperties prev_bp_level, T *data_cost_checkerboard_0, T *data_cost_checkerboard_1, T *data_cost_current_level, unsigned int offset_num, unsigned int bp_settings_disp_vals)
 Initialize the data costs at the "next" level up in the pyramid given that the data at the lower has been set. More...
 
template<RunData_t T, unsigned int DISP_VALS>
__global__ void InitializeMessageValsToDefaultKernel (beliefprop::BpLevelProperties current_bp_level, T *message_u_checkerboard_0, T *message_d_checkerboard_0, T *message_l_checkerboard_0, T *message_r_checkerboard_0, T *message_u_checkerboard_1, T *message_d_checkerboard_1, T *message_l_checkerboard_1, T *message_r_checkerboard_1, unsigned int bp_settings_disp_vals)
 Initialize the message values at each pixel of the current level to the default value. More...
 
template<RunData_t T, unsigned int DISP_VALS>
__global__ void RunBPIterationUsingCheckerboardUpdates (beliefprop::CheckerboardPart checkerboard_to_update, beliefprop::BpLevelProperties current_bp_level, T *data_cost_checkerboard_0, T *data_cost_checkerboard_1, T *message_u_checkerboard_0, T *message_d_checkerboard_0, T *message_l_checkerboard_0, T *message_r_checkerboard_0, T *message_u_checkerboard_1, T *message_d_checkerboard_1, T *message_l_checkerboard_1, T *message_r_checkerboard_1, float disc_k_bp, bool data_aligned, unsigned int bp_settings_disp_vals)
 Kernel function to run the current iteration of belief propagation in parallel using the checkerboard update method where half the pixels in the "checkerboard" scheme retrieve messages from each 4-connected neighbor and then update their message based on the retrieved messages and the data cost. More...
 
template<RunData_t T, unsigned int DISP_VALS>
__global__ void RunBPIterationUsingCheckerboardUpdates (beliefprop::CheckerboardPart checkerboard_to_update, beliefprop::BpLevelProperties current_bp_level, T *data_cost_checkerboard_0, T *data_cost_checkerboard_1, T *message_u_checkerboard_0, T *message_d_checkerboard_0, T *message_l_checkerboard_0, T *message_r_checkerboard_0, T *message_u_checkerboard_1, T *message_d_checkerboard_1, T *message_l_checkerboard_1, T *message_r_checkerboard_1, float disc_k_bp, bool data_aligned, unsigned int bp_settings_disp_vals, void *dst_processing)
 Kernel function to run the current iteration of belief propagation in parallel using the checkerboard update method where half the pixels in the "checkerboard" scheme retrieve messages from each 4-connected neighbor and then update their message based on the retrieved messages and the data cost. Function differs from counterpart overloaded function with same name in that it takes in allocated memory to use in processing. More...
 
template<RunData_t T, unsigned int DISP_VALS>
__global__ void CopyMsgDataToNextLevel (beliefprop::CheckerboardPart checkerboard_part, beliefprop::BpLevelProperties current_bp_level, beliefprop::BpLevelProperties next_bp_level, T *message_u_prev_checkerboard_0, T *message_d_prev_checkerboard_0, T *message_l_prev_checkerboard_0, T *message_r_prev_checkerboard_0, T *message_u_prev_checkerboard_1, T *message_d_prev_checkerboard_1, T *message_l_prev_checkerboard_1, T *message_r_prev_checkerboard_1, T *message_u_checkerboard_0, T *message_d_checkerboard_0, T *message_l_checkerboard_0, T *message_r_checkerboard_0, T *message_u_checkerboard_1, T *message_d_checkerboard_1, T *message_l_checkerboard_1, T *message_r_checkerboard_1, unsigned int bp_settings_disp_vals)
 Kernel to copy the computed BP message values at the current level to the corresponding locations at the "next" level down. The kernel works from the point of view of the pixel at the current level that is being copied to four corresponding pixels in the next level. More...
 
template<RunData_t T, unsigned int DISP_VALS>
__global__ void RetrieveOutputDisparity (beliefprop::BpLevelProperties current_bp_level, T *data_cost_checkerboard_0, T *data_cost_checkerboard_1, T *message_u_checkerboard_0, T *message_d_checkerboard_0, T *message_l_checkerboard_0, T *message_r_checkerboard_0, T *message_u_checkerboard_1, T *message_d_checkerboard_1, T *message_l_checkerboard_1, T *message_r_checkerboard_1, float *disparity_between_images_device, unsigned int bp_settings_disp_vals)
 Retrieve the best disparity estimate from image 1 to image 2 for each pixel in parallel. More...
 
__global__ void convertUnsignedIntImageToFloat (unsigned int *uint_image_pixels, float *float_image_pixels, unsigned int width_images, unsigned int height_images)
 
template<BpImData_t T>
__global__ void FilterImageAcross (const T *image_to_filter, float *filtered_image, unsigned int width_images, unsigned int height_images, const float *image_filter, unsigned int size_filter)
 Kernel to apply a horizontal filter on each pixel of the image in parallel the input image is stored as unsigned ints the output filtered image is returned as an array of floats. More...
 
template<BpImData_t T>
__global__ void FilterImageVertical (const T *image_to_filter, float *filtered_image, unsigned int width_images, unsigned int height_images, const float *image_filter, unsigned int size_filter)
 Kernel to apply a vertical filter on each pixel of the image in parallel the input image is stored as unsigned ints the output filtered image is returned as an array of floats. More...
 
__global__ void convertUnsignedIntImageToFloat (const unsigned int *uint_image_pixels, float *float_image_pixels, unsigned int width_images, unsigned int height_images)
 Kernel to convert the unsigned int pixels to float pixels in an image when smoothing is not desired but the pixels need to be converted to floats the input image is stored as unsigned ints output filtered image stored in float_image_pixels. More...
 

Detailed Description

Namespace to define global kernel functions for parallel belief propagation processing using CUDA.

Function Documentation

◆ convertUnsignedIntImageToFloat() [1/2]

__global__ void beliefprop_cuda::convertUnsignedIntImageToFloat ( const unsigned int *  uint_image_pixels,
float *  float_image_pixels,
unsigned int  width_images,
unsigned int  height_images 
)

Kernel to convert the unsigned int pixels to float pixels in an image when smoothing is not desired but the pixels need to be converted to floats the input image is stored as unsigned ints output filtered image stored in float_image_pixels.

Parameters
uint_image_pixels
float_image_pixels
width_images
height_images

◆ convertUnsignedIntImageToFloat() [2/2]

__global__ void beliefprop_cuda::convertUnsignedIntImageToFloat ( unsigned int *  uint_image_pixels,
float *  float_image_pixels,
unsigned int  width_images,
unsigned int  height_images 
)

Definition at line 43 of file KernelFilter.cu.

Here is the call graph for this function:

◆ CopyMsgDataToNextLevel()

template<RunData_t T, unsigned int DISP_VALS>
__global__ void beliefprop_cuda::CopyMsgDataToNextLevel ( beliefprop::CheckerboardPart  checkerboard_part,
beliefprop::BpLevelProperties  current_bp_level,
beliefprop::BpLevelProperties  next_bp_level,
T *  message_u_prev_checkerboard_0,
T *  message_d_prev_checkerboard_0,
T *  message_l_prev_checkerboard_0,
T *  message_r_prev_checkerboard_0,
T *  message_u_prev_checkerboard_1,
T *  message_d_prev_checkerboard_1,
T *  message_l_prev_checkerboard_1,
T *  message_r_prev_checkerboard_1,
T *  message_u_checkerboard_0,
T *  message_d_checkerboard_0,
T *  message_l_checkerboard_0,
T *  message_r_checkerboard_0,
T *  message_u_checkerboard_1,
T *  message_d_checkerboard_1,
T *  message_l_checkerboard_1,
T *  message_r_checkerboard_1,
unsigned int  bp_settings_disp_vals 
)

Kernel to copy the computed BP message values at the current level to the corresponding locations at the "next" level down. The kernel works from the point of view of the pixel at the current level that is being copied to four corresponding pixels in the next level.

Template Parameters
T
DISP_VALS
Parameters
checkerboard_part
current_bp_level
next_bp_level
message_u_prev_checkerboard_0
message_d_prev_checkerboard_0
message_l_prev_checkerboard_0
message_r_prev_checkerboard_0
message_u_prev_checkerboard_1
message_d_prev_checkerboard_1
message_l_prev_checkerboard_1
message_r_prev_checkerboard_1
message_u_checkerboard_0
message_d_checkerboard_0
message_l_checkerboard_0
message_r_checkerboard_0
message_u_checkerboard_1
message_d_checkerboard_1
message_l_checkerboard_1
message_r_checkerboard_1
bp_settings_disp_vals

Definition at line 310 of file KernelBpStereo.cu.

Here is the call graph for this function:

◆ FilterImageAcross()

template<BpImData_t T>
__global__ void beliefprop_cuda::FilterImageAcross ( const T *  image_to_filter,
float *  filtered_image,
unsigned int  width_images,
unsigned int  height_images,
const float *  image_filter,
unsigned int  size_filter 
)

Kernel to apply a horizontal filter on each pixel of the image in parallel the input image is stored as unsigned ints the output filtered image is returned as an array of floats.

Template Parameters
T
Parameters
image_to_filter
filtered_image
width_images
height_images
image_filter
size_filter

Definition at line 62 of file KernelFilter.cu.

Here is the call graph for this function:

◆ FilterImageVertical()

template<BpImData_t T>
__global__ void beliefprop_cuda::FilterImageVertical ( const T *  image_to_filter,
float *  filtered_image,
unsigned int  width_images,
unsigned int  height_images,
const float *  image_filter,
unsigned int  size_filter 
)

Kernel to apply a vertical filter on each pixel of the image in parallel the input image is stored as unsigned ints the output filtered image is returned as an array of floats.

Template Parameters
T
Parameters
image_to_filter
filtered_image
width_images
height_images
image_filter
size_filter

Definition at line 82 of file KernelFilter.cu.

Here is the call graph for this function:

◆ InitializeBottomLevelData()

template<RunData_t T, unsigned int DISP_VALS>
__global__ void beliefprop_cuda::InitializeBottomLevelData ( beliefprop::BpLevelProperties  current_bp_level,
float *  image_1_pixels_device,
float *  image_2_pixels_device,
T *  data_cost_stereo_checkerboard_0,
T *  data_cost_stereo_checkerboard_1,
float  lambda_bp,
float  data_k_bp,
unsigned int  bp_settings_disp_vals 
)

Initialize the "data cost" for each possible disparity between the two full-sized input images ("bottom" of the image pyramid). The image data is stored in the image_1_pixels_device and image_2_pixels_device arrays.

Template Parameters
T
DISP_VALS
Parameters
current_bp_level
image_1_pixels_device
image_2_pixels_device
data_cost_stereo_checkerboard_0
data_cost_stereo_checkerboard_1
lambda_bp
data_k_bp
bp_settings_disp_vals

Definition at line 68 of file KernelBpStereo.cu.

Here is the call graph for this function:

◆ InitializeCurrentLevelData()

template<RunData_t T, unsigned int DISP_VALS>
__global__ void beliefprop_cuda::InitializeCurrentLevelData ( beliefprop::CheckerboardPart  checkerboard_part,
beliefprop::BpLevelProperties  current_bp_level,
beliefprop::BpLevelProperties  prev_bp_level,
T *  data_cost_checkerboard_0,
T *  data_cost_checkerboard_1,
T *  data_cost_current_level,
unsigned int  offset_num,
unsigned int  bp_settings_disp_vals 
)

Initialize the data costs at the "next" level up in the pyramid given that the data at the lower has been set.

Template Parameters
T
DISP_VALS
Parameters
checkerboard_part
current_bp_level
prev_bp_level
data_cost_checkerboard_0
data_cost_checkerboard_1
data_cost_current_level
offset_num
bp_settings_disp_vals

Definition at line 108 of file KernelBpStereo.cu.

Here is the call graph for this function:

◆ InitializeMessageValsToDefaultKernel()

template<RunData_t T, unsigned int DISP_VALS>
__global__ void beliefprop_cuda::InitializeMessageValsToDefaultKernel ( beliefprop::BpLevelProperties  current_bp_level,
T *  message_u_checkerboard_0,
T *  message_d_checkerboard_0,
T *  message_l_checkerboard_0,
T *  message_r_checkerboard_0,
T *  message_u_checkerboard_1,
T *  message_d_checkerboard_1,
T *  message_l_checkerboard_1,
T *  message_r_checkerboard_1,
unsigned int  bp_settings_disp_vals 
)

Initialize the message values at each pixel of the current level to the default value.

Template Parameters
T
DISP_VALS
Parameters
current_bp_level
message_u_checkerboard_0
message_d_checkerboard_0
message_l_checkerboard_0
message_r_checkerboard_0
message_u_checkerboard_1
message_d_checkerboard_1
message_l_checkerboard_1
message_r_checkerboard_1
bp_settings_disp_vals

Definition at line 146 of file KernelBpStereo.cu.

Here is the call graph for this function:

◆ RetrieveOutputDisparity()

template<RunData_t T, unsigned int DISP_VALS>
__global__ void beliefprop_cuda::RetrieveOutputDisparity ( beliefprop::BpLevelProperties  current_bp_level,
T *  data_cost_checkerboard_0,
T *  data_cost_checkerboard_1,
T *  message_u_checkerboard_0,
T *  message_d_checkerboard_0,
T *  message_l_checkerboard_0,
T *  message_r_checkerboard_0,
T *  message_u_checkerboard_1,
T *  message_d_checkerboard_1,
T *  message_l_checkerboard_1,
T *  message_r_checkerboard_1,
float *  disparity_between_images_device,
unsigned int  bp_settings_disp_vals 
)

Retrieve the best disparity estimate from image 1 to image 2 for each pixel in parallel.

Template Parameters
T
DISP_VALS
Parameters
current_bp_level
data_cost_checkerboard_0
data_cost_checkerboard_1
message_u_checkerboard_0
message_d_checkerboard_0
message_l_checkerboard_0
message_r_checkerboard_0
message_u_checkerboard_1
message_d_checkerboard_1
message_l_checkerboard_1
message_r_checkerboard_1
disparity_between_images_device
bp_settings_disp_vals

Definition at line 365 of file KernelBpStereo.cu.

Here is the call graph for this function:

◆ RunBPIterationUsingCheckerboardUpdates() [1/2]

template<RunData_t T, unsigned int DISP_VALS>
__global__ void beliefprop_cuda::RunBPIterationUsingCheckerboardUpdates ( beliefprop::CheckerboardPart  checkerboard_to_update,
beliefprop::BpLevelProperties  current_bp_level,
T *  data_cost_checkerboard_0,
T *  data_cost_checkerboard_1,
T *  message_u_checkerboard_0,
T *  message_d_checkerboard_0,
T *  message_l_checkerboard_0,
T *  message_r_checkerboard_0,
T *  message_u_checkerboard_1,
T *  message_d_checkerboard_1,
T *  message_l_checkerboard_1,
T *  message_r_checkerboard_1,
float  disc_k_bp,
bool  data_aligned,
unsigned int  bp_settings_disp_vals 
)

Kernel function to run the current iteration of belief propagation in parallel using the checkerboard update method where half the pixels in the "checkerboard" scheme retrieve messages from each 4-connected neighbor and then update their message based on the retrieved messages and the data cost.

Template Parameters
T
DISP_VALS
Parameters
checkerboard_to_update
current_bp_level
data_cost_checkerboard_0
data_cost_checkerboard_1
message_u_checkerboard_0
message_d_checkerboard_0
message_l_checkerboard_0
message_r_checkerboard_0
message_u_checkerboard_1
message_d_checkerboard_1
message_l_checkerboard_1
message_r_checkerboard_1
disc_k_bp
data_aligned
bp_settings_disp_vals

Definition at line 197 of file KernelBpStereo.cu.

Here is the call graph for this function:

◆ RunBPIterationUsingCheckerboardUpdates() [2/2]

template<RunData_t T, unsigned int DISP_VALS>
__global__ void beliefprop_cuda::RunBPIterationUsingCheckerboardUpdates ( beliefprop::CheckerboardPart  checkerboard_to_update,
beliefprop::BpLevelProperties  current_bp_level,
T *  data_cost_checkerboard_0,
T *  data_cost_checkerboard_1,
T *  message_u_checkerboard_0,
T *  message_d_checkerboard_0,
T *  message_l_checkerboard_0,
T *  message_r_checkerboard_0,
T *  message_u_checkerboard_1,
T *  message_d_checkerboard_1,
T *  message_l_checkerboard_1,
T *  message_r_checkerboard_1,
float  disc_k_bp,
bool  data_aligned,
unsigned int  bp_settings_disp_vals,
void *  dst_processing 
)

Kernel function to run the current iteration of belief propagation in parallel using the checkerboard update method where half the pixels in the "checkerboard" scheme retrieve messages from each 4-connected neighbor and then update their message based on the retrieved messages and the data cost. Function differs from counterpart overloaded function with same name in that it takes in allocated memory to use in processing.

Template Parameters
T
DISP_VALS
Parameters
checkerboard_to_update
current_bp_level
data_cost_checkerboard_0
data_cost_checkerboard_1
message_u_checkerboard_0
message_d_checkerboard_0
message_l_checkerboard_0
message_r_checkerboard_0
message_u_checkerboard_1
message_d_checkerboard_1
message_l_checkerboard_1
message_r_checkerboard_1
disc_k_bp
data_aligned
bp_settings_disp_vals
dst_processing

Definition at line 252 of file KernelBpStereo.cu.

Here is the call graph for this function: