Help save net neutrality! Learn more.
Close

[289b89]: / contrib / brl / bseg / boxm2 / ocl / cl / bit / compute_index.cl  Maximize  Restore  History

Download this file

139 lines (118 with data), 5.1 kB

  1
  2
  3
  4
  5
  6
  7
  8
  9
 10
 11
 12
 13
 14
 15
 16
 17
 18
 19
 20
 21
 22
 23
 24
 25
 26
 27
 28
 29
 30
 31
 32
 33
 34
 35
 36
 37
 38
 39
 40
 41
 42
 43
 44
 45
 46
 47
 48
 49
 50
 51
 52
 53
 54
 55
 56
 57
 58
 59
 60
 61
 62
 63
 64
 65
 66
 67
 68
 69
 70
 71
 72
 73
 74
 75
 76
 77
 78
 79
 80
 81
 82
 83
 84
 85
 86
 87
 88
 89
 90
 91
 92
 93
 94
 95
 96
 97
 98
 99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
115
116
117
118
119
120
121
122
123
124
125
126
127
128
129
130
131
132
133
134
135
136
137
#pragma OPENCL EXTENSION cl_khr_global_int32_base_atomics: enable
#if NVIDIA
#pragma OPENCL EXTENSION cl_khr_gl_sharing : enable
#endif
#ifdef COMPINDEX
//need to define a struct of type AuxArgs with auxiliary arguments
// to supplement cast ray args
typedef struct
{
__global float* alpha;
float* expdepth;
float* expdepthsqr;
float* probsum;
float* t;
float* vis;
} AuxArgs;
//forward declare cast ray (so you can use it)
/*void cast_ray(int,int,float,float,float,float,float,float,
__constant RenderSceneInfo*, __global int4*,
__local uchar16*, __constant uchar *,__local uchar *,
float*, AuxArgs);*/
void cast_ray_render_vis(int,int,float,float,float,float,float,float,
__constant RenderSceneInfo*, __global int4*,
__local uchar16*, __constant uchar *,__local uchar *,
float*, AuxArgs);
void step_cell_compute_index(float depth,
float block_len,
__global float * alpha_data,
int data_ptr,
float d,
float * vis,
float * expected_depth,
float * expected_depth_square,
float * probsum,
float * t)
{
float alpha = alpha_data[data_ptr];
float diff_omega=exp(-alpha*d);
float omega=(*vis) * (1.0f - diff_omega);
(*probsum)+=omega;
(*vis) *= diff_omega;
(*expected_depth)+=depth*omega;
(*expected_depth_square)+=depth*depth*omega;
(*t)=depth*block_len;
}
__kernel
void
compute_loc_index(
__constant RenderSceneInfo * linfo,
__global float4 * directions,
__constant int * ray_size,
__global int4 * tree_array, // tree structure for each block
__constant uchar * bit_lookup, // used to get data_index
__global float * alpha_array,
__constant float4 * ray_o,
__constant float * max_dist,
__global float * exp_depth_buf,
__global float * vis_buf,
__global float * prob_buf,
__global float * t_infinity_buf,
__local uchar16 * local_tree, // cache current tree into local memory
__local uchar * cumsum
)
{
int gid=get_global_id(0);
//exp_depth_buf[ gid ] = -100.0f;
//vis_buf[ gid ] = gid+1;
//prob_buf[ 0 ] = 1000.0f;
if (gid >= *ray_size)
return;
float4 ray_d = directions[ gid ];
//declare ray
float ray_ox, ray_oy, ray_oz, ray_dx, ray_dy, ray_dz;
calc_scene_ray_generic_cam(linfo, *ray_o, ray_d, &ray_ox, &ray_oy, &ray_oz, &ray_dx, &ray_dy, &ray_dz);
float expdepth = 0.0f;
float expdepthsqr = 0.0f;
float probsum = prob_buf[ gid ];
float vis_rec = vis_buf[ gid ];
float t = t_infinity_buf[ gid ];
float tfar_max = (*max_dist)/linfo->block_len;
AuxArgs aux_args;
aux_args.alpha = alpha_array;
aux_args.expdepth = &expdepth;
aux_args.expdepthsqr = &expdepthsqr;
aux_args.probsum = &probsum;
aux_args.t = &t;
aux_args.vis = &vis_rec;
/*float vis = 1.0;
cast_ray( 1, 1,
ray_ox, ray_oy, ray_oz,
ray_dx, ray_dy, ray_dz,
linfo, tree_array, //scene info
local_tree, bit_lookup, cumsum, &vis, aux_args); //utility info*/
cast_ray_render_vis( 1, 1,
ray_ox, ray_oy, ray_oz,
ray_dx, ray_dy, ray_dz,
linfo, tree_array, //scene info
local_tree, bit_lookup, cumsum, &tfar_max, aux_args,0,MAXFLOAT); //utility info
//store values at the end of this block
exp_depth_buf[ gid ] += (* aux_args.expdepth)*linfo->block_len;
prob_buf[ gid ] = (* aux_args.probsum);
vis_buf[ gid ] = vis_rec;
t_infinity_buf[ gid ] = (* aux_args.t);
}
__kernel void normalize_index_depth_kernel(
__global float * exp_depth_buf,
__global float * prob_buf,
__global float * t_infinity_buf,
__global float * sub_block_dim)
{
int gid=get_global_id(0);
//normalize
float prob = prob_buf[gid];
float mean = exp_depth_buf[gid] + t_infinity_buf[gid]*prob;
//float mean = exp_depth_buf[gid] + t_infinity_buf[gid]*prob * (*sub_block_dim);
exp_depth_buf[gid]=mean;
}
#endif //COMPINDEX