This file defines the template specialization to perform belief propagation using half precision for disparity map estimation from stereo images on CUDA to prevent overflow in val_to_normalize message value computation shouldn't be an issue if using bfloat instead of half so recommend using that instead of using these template specialization functions.
More...
|
| template<unsigned int DISP_VALS> |
| __device__ void | MsgStereoHalf (unsigned int x_val, unsigned int y_val, const BpLevel< T > ¤t_bp_level, half messages_neighbor_1[DISP_VALS], half messages_neighbor_2[DISP_VALS], half messages_neighbor_3[DISP_VALS], half data_costs[DISP_VALS], half *dst_message_array, half disc_k_bp, bool data_aligned) |
| |
| template<beliefprop::MessageComp M> |
| __device__ void | MsgStereoHalf (unsigned int x_val, unsigned int y_val, const BpLevel< T > ¤t_bp_level, half *prev_u_messageArray, half *prev_d_messageArray, half *prev_l_messageArray, half *prev_r_messageArray, half *data_message_array, half *dst_message_array, half disc_k_bp, bool data_aligned, unsigned int bp_settings_disp_vals, half *dst_processing, unsigned int checkerboard_adjustment, unsigned int offset_data) |
| |
| template<> |
| __device__ void | MsgStereo< half, half, beliefprop::MessageComp::kUMessage > (unsigned int x_val, unsigned int y_val, const BpLevel< T > ¤t_bp_level, half *prev_u_messageArray, half *prev_d_messageArray, half *prev_l_messageArray, half *prev_r_messageArray, half *data_message_array, half *dst_message_array, half disc_k_bp, bool data_aligned, unsigned int bp_settings_disp_vals, half *dst_processing, unsigned int checkerboard_adjustment, unsigned int offset_data) |
| |
| template<> |
| __device__ void | MsgStereo< half, half, beliefprop::MessageComp::kDMessage > (unsigned int x_val, unsigned int y_val, const BpLevel< T > ¤t_bp_level, half *prev_u_messageArray, half *prev_d_messageArray, half *prev_l_messageArray, half *prev_r_messageArray, half *data_message_array, half *dst_message_array, half disc_k_bp, bool data_aligned, unsigned int bp_settings_disp_vals, half *dst_processing, unsigned int checkerboard_adjustment, unsigned int offset_data) |
| |
| template<> |
| __device__ void | MsgStereo< half, half, beliefprop::MessageComp::kLMessage > (unsigned int x_val, unsigned int y_val, const BpLevel< T > ¤t_bp_level, half *prev_u_messageArray, half *prev_d_messageArray, half *prev_l_messageArray, half *prev_r_messageArray, half *data_message_array, half *dst_message_array, half disc_k_bp, bool data_aligned, unsigned int bp_settings_disp_vals, half *dst_processing, unsigned int checkerboard_adjustment, unsigned int offset_data) |
| |
| template<> |
| __device__ void | MsgStereo< half, half, beliefprop::MessageComp::kRMessage > (unsigned int x_val, unsigned int y_val, const BpLevel< T > ¤t_bp_level, half *prev_u_messageArray, half *prev_d_messageArray, half *prev_l_messageArray, half *prev_r_messageArray, half *data_message_array, half *dst_message_array, half disc_k_bp, bool data_aligned, unsigned int bp_settings_disp_vals, half *dst_processing, unsigned int checkerboard_adjustment, unsigned int offset_data) |
| |
| template<> |
| __device__ void | MsgStereo< half, half, kDispVals0 > (unsigned int x_val, unsigned int y_val, const BpLevel< T > ¤t_bp_level, half messages_neighbor_1[kDispVals0], half messages_neighbor_2[kDispVals0], half messages_neighbor_3[kDispVals0], half data_costs[kDispVals0], half *dst_message_array, half disc_k_bp, bool data_aligned) |
| |
| template<> |
| __device__ void | MsgStereo< half, half, kDispVals1 > (unsigned int x_val, unsigned int y_val, const BpLevel< T > ¤t_bp_level, half messages_neighbor_1[kDispVals1], half messages_neighbor_2[kDispVals1], half messages_neighbor_3[kDispVals1], half data_costs[kDispVals1], half *dst_message_array, half disc_k_bp, bool data_aligned) |
| |
| template<> |
| __device__ void | MsgStereo< half, half, kDispVals2 > (unsigned int x_val, unsigned int y_val, const BpLevel< T > ¤t_bp_level, half messages_neighbor_1[kDispVals2], half messages_neighbor_2[kDispVals2], half messages_neighbor_3[kDispVals2], half data_costs[kDispVals2], half *dst_message_array, half disc_k_bp, bool data_aligned) |
| |
| template<> |
| __device__ void | MsgStereo< half, half, kDispVals3 > (unsigned int x_val, unsigned int y_val, const BpLevel< T > ¤t_bp_level, half messages_neighbor_1[kDispVals3], half messages_neighbor_2[kDispVals3], half messages_neighbor_3[kDispVals3], half data_costs[kDispVals3], half *dst_message_array, half disc_k_bp, bool data_aligned) |
| |
| template<> |
| __device__ void | MsgStereo< half, half, kDispVals4 > (unsigned int x_val, unsigned int y_val, const BpLevel< T > ¤t_bp_level, half messages_neighbor_1[kDispVals4], half messages_neighbor_2[kDispVals4], half messages_neighbor_3[kDispVals4], half data_costs[kDispVals4], half *dst_message_array, half disc_k_bp, bool data_aligned) |
| |
| template<> |
| __device__ void | MsgStereo< half, half, kDispVals5 > (unsigned int x_val, unsigned int y_val, const BpLevel< T > ¤t_bp_level, half messages_neighbor_1[kDispVals5], half messages_neighbor_2[kDispVals5], half messages_neighbor_3[kDispVals5], half data_costs[kDispVals5], half *dst_message_array, half disc_k_bp, bool data_aligned) |
| |
| template<> |
| __device__ void | MsgStereo< half, half, kDispVals6 > (unsigned int x_val, unsigned int y_val, const BpLevel< T > ¤t_bp_level, half messages_neighbor_1[kDispVals6], half messages_neighbor_2[kDispVals6], half messages_neighbor_3[kDispVals6], half data_costs[kDispVals6], half *dst_message_array, half disc_k_bp, bool data_aligned) |
| |
This file defines the template specialization to perform belief propagation using half precision for disparity map estimation from stereo images on CUDA to prevent overflow in val_to_normalize message value computation shouldn't be an issue if using bfloat instead of half so recommend using that instead of using these template specialization functions.
- Author
- Scott Grauer-Gray
- Copyright
- Copyright (c) 2024
Definition in file KernelBpStereoHalf.cu.
