1//THIS IS UPDATE BIT SCENE OPT
2//Created Sept 30, 2010,
3//Implements the parallel work group segmentation algorithm.
4#pragma OPENCL EXTENSION cl_khr_global_int32_base_atomics: enable
5#if NVIDIA
6 #pragma OPENCL EXTENSION cl_khr_gl_sharing : enable
7#endif
8
9#ifdef SEGLENVIS
10
11typedef struct
12{
13  __global float * alpha;
14  __global int  * seg_len;
15  __global int  * vis_array;
16  __local  short2* ray_bundle_array;
17  __local  int*    cell_ptrs;
18 // __local  float*  cached_vis;
19           float*  ray_vis;
20           float*  last_vis;
21} AuxArgs;
22
23
24//bayes step cell functor
25void step_cell_seglen_vis(AuxArgs aux_args, int data_ptr, uchar llid, float d)
26{
27    //slow beta calculation ----------------------------------------------------
28
29    int seg_int = convert_int_rte(d * SEGLEN_FACTOR);
30    atom_add(&aux_args.seg_len[data_ptr], seg_int);
31    //calculate this ray's contribution to beta
32    float alpha = aux_args.alpha[data_ptr];
33    //float cell_vis = (* aux_args.ray_vis) * d;
34    float cell_vis = *(aux_args.last_vis);
35
36    //update ray_pre and ray_vis
37    float temp = exp(-alpha * d);
38    // updated visibility probability
39    (* aux_args.ray_vis) *= temp;
40
41    // keep track of vis at last "empty" cell to prevent self-shadowing within uncertain regions
42    constexpr float passthrough_prob_thresh = 0.9;
43    if (temp >= passthrough_prob_thresh) {
44       *(aux_args.last_vis) = *(aux_args.ray_vis);
45    }
46
47    //discretize and store beta and vis contribution
48    int vis_int  = convert_int_rte(d * cell_vis * SEGLEN_FACTOR);
49    atom_add(&aux_args.vis_array[data_ptr], vis_int);
50    //--------------------------------------------------------------------------
51}
52//forward declare cast ray (so you can use it)
53void cast_ray(int,int,float,float,float,float,float,float,__constant RenderSceneInfo*,
54              __global int4*,local uchar16*,constant uchar *,local uchar *,float*,AuxArgs,float tnear, float tfar);
55
56__kernel
57void
58seg_len_and_vis_main(__constant  RenderSceneInfo    * linfo,
59                     __global    int4               * tree_array,        // tree structure for each block
60                     __global    float              * alpha_array,       // alpha for each block
61                     __global    int                * aux_array0,        // four aux arrays strung together
62                     __global    int                * aux_array1,        // four aux arrays strung together
63                     __constant  uchar              * bit_lookup,        // used to get data_index
64                     __global    float4             * ray_origins,
65                     __global    float4             * ray_directions,
66                     __global    uint4              * imgdims,           // dimensions of the input image
67                     __global    float              * vis_image,         // visibility image (for keeping vis across blocks)
68                     __global    float              * last_vis_image,    // like visibility image, but updates only in "empty" regions to prevent self-shadowing within uncertain regions
69                     __global    float              * output,
70                     __local     uchar16            * local_tree,        // cache current tree into local memory
71                     __local     short2             * ray_bundle_array,  // gives information for which ray takes over in the workgroup
72                     __local     int                * cell_ptrs,         // local list of cell_ptrs (cells that are hit by this workgroup
73                     __local     float              * cached_vis,        // cached vis used to sum up vis contribution locally
74                     __local     uchar              * cumsum)            // cumulative sum for calculating data pointer
75{
76  // get local id (0-63 for an 8x8) of this patch
77  uchar llid = (uchar)(get_local_id(0) + get_local_size(0)*get_local_id(1));
78
79  // initialize pre-broken ray information (non broken rays will be re initialized)
80  ray_bundle_array[llid] = (short2) (-1, 0);
81  cell_ptrs[llid] = -1;
82
83  //----------------------------------------------------------------------------
84  // get image coordinates and camera,
85  // check for validity before proceeding
86  //----------------------------------------------------------------------------
87  int i=0,j=0;
88  i=get_global_id(0);
89  j=get_global_id(1);
90
91  // check to see if the thread corresponds to an actual pixel as in some
92  // cases #of threads will be more than the pixels.
93  if (i>=(*imgdims).z || j>=(*imgdims).w || i<(*imgdims).x || j<(*imgdims).y)
94    return;
95  float vis = vis_image[j*get_global_size(0) + i];
96  float last_vis = last_vis_image[j*get_global_size(0) + i];
97  float vis0 = 1.0; // don't want to terminate raytrace early, even if visibility is 0
98
99  barrier(CLK_LOCAL_MEM_FENCE);
100
101  //----------------------------------------------------------------------------
102  // we know i,j map to a point on the image,
103  // BEGIN RAY TRACE
104  //----------------------------------------------------------------------------
105  float4 ray_o = ray_origins[ j*get_global_size(0) + i ];
106  float4 ray_d = ray_directions[ j*get_global_size(0) + i ];
107  float ray_ox, ray_oy, ray_oz, ray_dx, ray_dy, ray_dz;
108
109  calc_scene_ray_generic_cam(linfo, ray_o, ray_d, &ray_ox, &ray_oy, &ray_oz, &ray_dx, &ray_dy, &ray_dz);
110
111  //----------------------------------------------------------------------------
112  // we know i,j map to a point on the image, have calculated ray
113  // BEGIN RAY TRACE
114  //----------------------------------------------------------------------------
115  AuxArgs aux_args;
116  aux_args.alpha       = alpha_array;
117  aux_args.seg_len     = aux_array0;
118  aux_args.vis_array   = aux_array1;
119
120  aux_args.ray_bundle_array = ray_bundle_array;
121  aux_args.cell_ptrs = cell_ptrs;
122  //aux_args.cached_vis = cached_vis;
123  aux_args.ray_vis = &vis;
124  aux_args.last_vis = &last_vis;
125  cast_ray( i, j,
126            ray_ox, ray_oy, ray_oz,
127            ray_dx, ray_dy, ray_dz,
128            linfo, tree_array,                                  //scene info
129            local_tree, bit_lookup, cumsum, &vis0, aux_args,0,MAXFLOAT);    //utility info
130
131  //write out vis and pre
132  vis_image[j*get_global_size(0)+i] = vis;
133  last_vis_image[j*get_global_size(0)+i] = last_vis;
134
135}
136#endif
137
138#ifdef UPDATE_SUN_VIS
139__kernel void
140update_visibilities_main(__constant  RenderSceneInfo    * linfo,
141                         __global float* aux_array0,
142                         __global float* aux_array1,
143                         __global float* aux_output)
144{
145  int gid=get_global_id(0);
146  int datasize = linfo->data_len ;
147  if (gid<datasize)
148  {
149    int obs0= as_int(aux_array0[gid]);
150    int obs1= as_int(aux_array1[gid]);
151    if(obs0>0)
152        aux_output[gid] = ((float)obs1)/((float)obs0);
153    else
154        aux_output[gid] = 0.0;
155  }
156
157}
158#endif
159