hier_bp.cl File Reference

Macros
#define	HIER_ITER   4
#define	DISP_RANGE   64
#define	DISP_CAMERA   32
#define	LOCAL_SIZE   8

Functions
__kernel void	initializeHierCost (__global float *datacost, __global float *hierCost, int width, int height, int dispRange)
void	gpu_comp_msg (__global float *m1, __global float *m2, __global float *m3, __global float *data, __global float *dest, int MAX_DISP, float DISCONT_COST, float CONT_COST)
void	gpu_comp_msg_local (__local float *m1, __local float *m2, __local float *m3, float *data, float *dest, int MAX_DISP, float DISCONT_COST, float CONT_COST)
__kernel void	hierBP (__global float *temp_m_u, __global float *temp_m_d, __global float *temp_m_l, __global float *temp_m_r, __global float *datacost, __global float *m_u, __global float *m_d, __global float *m_l, __global float *m_r, __global float *gpu_zero, int width, int height, int dispRange, int dist, float DISCONT_COST, float CONT_COST)
__kernel void	hierBP_local (__global float *temp_m_u, __global float *temp_m_d, __global float *temp_m_l, __global float *temp_m_r, __global float *datacost, __global float *m_u, __global float *m_d, __global float *m_l, __global float *m_r, __global float *gpu_zero, int width, int height, int dispRange, int dist, float DISCONT_COST, float CONT_COST)
__kernel void	updateCostLayer (__global float *m_u, __global float *m_d, __global float *m_l, __global float *m_r, int width, int height, int dispRange, int dist)
__kernel void	finalBP (__global float *m_u, __global float *m_d, __global float *m_l, __global float *m_r, __global float *datacost, __global float *gpu_zero, int width, int height, int dispRange)

Macro Definition Documentation

DISP_CAMERA

#define DISP_CAMERA   32

DISP_RANGE

#define DISP_RANGE   64

HIER_ITER

#define HIER_ITER   4

LOCAL_SIZE

#define LOCAL_SIZE   8

Function Documentation

finalBP()

__kernel void finalBP	(	__global float *	m_u,
		__global float *	m_d,
		__global float *	m_l,
		__global float *	m_r,
		__global float *	datacost,
		__global float *	gpu_zero,
		int	width,
		int	height,
		int	dispRange
	)

The final belief propagation in the base layer

Parameters

m_u	The array storage for node in up direction
m_d	The array storage for node in down direction
m_l	The array storage for node in left direction
m_r	The array storage for node in right direction
datacost	The 3D datacost
gpu_zero	Array containing zero to be used in GPU
width	The image width
height	The image height
dispRange	The disparity search range

gpu_comp_msg()

void gpu_comp_msg	(	__global float *	m1,
		__global float *	m2,
		__global float *	m3,
		__global float *	data,
		__global float *	dest,
		int	MAX_DISP,
		float	DISCONT_COST,
		float	CONT_COST
	)

Belief message propagation computation. This uses only global memory

Parameters

m1	The 1st neighbor node
m2	The 2nd neighbor node
m3	The 3rd neighbor node
data	The current node value
dest	The node to store the compute belief message
MAX_DISP	The disparity search range
DISCONT_COST	The edge hyperparameter
CONT_COST	The smoothing hyperparameter

gpu_comp_msg_local()

void gpu_comp_msg_local	(	__local float *	m1,
		__local float *	m2,
		__local float *	m3,
		float *	data,
		float *	dest,
		int	MAX_DISP,
		float	DISCONT_COST,
		float	CONT_COST
	)

Belief message propagation computation. This uses local memory (faster) if the GPU supports

Parameters

m1	The 1st neighbor node
m2	The 2nd neighbor node
m3	The 3rd neighbor node
data	The current node value
dest	The node to store the compute belief message
MAX_DISP	The disparity search range
DISCONT_COST	The edge hyperparameter
CONT_COST	The smoothing hyperparameter

hierBP()

__kernel void hierBP	(	__global float *	temp_m_u,
		__global float *	temp_m_d,
		__global float *	temp_m_l,
		__global float *	temp_m_r,
		__global float *	datacost,
		__global float *	m_u,
		__global float *	m_d,
		__global float *	m_l,
		__global float *	m_r,
		__global float *	gpu_zero,
		int	width,
		int	height,
		int	dispRange,
		int	dist,
		float	DISCONT_COST,
		float	CONT_COST
	)

The interface for hierarchal belief propagation using only global memory

Parameters

temp_m_u	The temporary array storage for node in up direction
temp_m_d	The temporary array storage for node in down direction
temp_m_l	The temporary array storage for node in left direction
temp_m_r	The temporary array storage for node in right direction
datacost	The 3D datacost
m_u	The array storage for node in up direction
m_d	The array storage for node in down direction
m_l	The array storage for node in left direction
m_r	The array storage for node in right direction
gpu_zero	Array containing zero to be used in GPU
width	The image width
height	The image height
dispRange	The disparity search range
dist	The distance value used for indexing
DISCONT_COST	The edge hyperparameter
CONT_COST	The smoothing hyperparameter

hierBP_local()

__kernel void hierBP_local	(	__global float *	temp_m_u,
		__global float *	temp_m_d,
		__global float *	temp_m_l,
		__global float *	temp_m_r,
		__global float *	datacost,
		__global float *	m_u,
		__global float *	m_d,
		__global float *	m_l,
		__global float *	m_r,
		__global float *	gpu_zero,
		int	width,
		int	height,
		int	dispRange,
		int	dist,
		float	DISCONT_COST,
		float	CONT_COST
	)

The interface for hierarchal belief propagation using local memory

Parameters

temp_m_u	The temporary array storage for node in up direction
temp_m_d	The temporary array storage for node in down direction
temp_m_l	The temporary array storage for node in left direction
temp_m_r	The temporary array storage for node in right direction
datacost	The 3D datacost
m_u	The array storage for node in up direction
m_d	The array storage for node in down direction
m_l	The array storage for node in left direction
m_r	The array storage for node in right direction
gpu_zero	Array containing zero to be used in GPU
width	The image width
height	The image height
dispRange	The disparity search range
dist	The distance value used for indexing
DISCONT_COST	The edge hyperparameter
CONT_COST	The smoothing hyperparameter

initializeHierCost()

__kernel void initializeHierCost	(	__global float *	datacost,
		__global float *	hierCost,
		int	width,
		int	height,
		int	dispRange
	)

Interpolate the initial cost (base) into a hierachal structure datacost

Parameters

datacost	The initial 3D cost
hierCost	Hierarchal 3D cost
width	The image width
height	The image height
dispRange	The stereo disparity search range

updateCostLayer()

__kernel void updateCostLayer	(	__global float *	m_u,
		__global float *	m_d,
		__global float *	m_l,
		__global float *	m_r,
		int	width,
		int	height,
		int	dispRange,
		int	dist
	)

Update the computed cost into the next hierarchal layer

Parameters

m_u	The array storage for node in up direction
m_d	The array storage for node in down direction
m_l	The array storage for node in left direction
m_r	The array storage for node in right direction
width	The image width
height	The image height
dispRange	The disparity search range
dist	The distance value used for indexing