1#if NVIDIA
2 #pragma OPENCL EXTENSION cl_khr_gl_sharing : enable
3#endif
4
5#ifdef RENDER
6//need to define a struct of type AuxArgs with auxiliary arguments
7// to supplement cast ray args
8// NOTE THAT RGB MOG is stored as follows:
9// [mu_0R, mu_0G, mu_0B, x] [sig_0R, sig_0G, sig_0B, w_0] [mu_1R, mu_1G, mu_1B, x] [sig_1R, sig_1G, sig_1B, x];
10typedef struct
11{
12  __global float*   alpha; //alpha
13  __global int2 *   mog;    //mixture of 2 3-d gaussians (as uchar 16
14           float*   vis;
15           float4*  expint; //expected intensity (3 channels)
16} AuxArgs;
17
18//forward declare cast ray (so you can use it)
19void cast_ray(int,int,float,float,float,float,float,float,__constant RenderSceneInfo*,
20              __global int4*,__local uchar16*,__constant uchar*,__local uchar*,float*,AuxArgs, float tnear, float tfar);
21__kernel
22void
23render_bit_scene( __constant  RenderSceneInfo    * linfo,
24                  __global    int4               * tree_array,
25                  __global    float              * alpha_array,
26                  __global    int2               * mixture_array,
27                  __global    float4             * ray_origins,
28                  __global    float4             * ray_directions,
29                  __global    float              * nearfarplanes,
30                  __global    float4             * exp_image,      // input image and store vis_inf and pre_inf
31                  __global    uint4              * exp_image_dims,
32                  __global    float              * output,
33                  __constant  uchar              * bit_lookup,
34                  __global    float              * vis_image,
35                  __global    float              * max_omega_image,
36                  __local     uchar16            * local_tree,
37                  __local     uchar              * cumsum,        //cumulative sum helper for data pointer
38                  __local     int                * imIndex)
39{
40  //----------------------------------------------------------------------------
41  //get local id (0-63 for an 8x8) of this patch + image coordinates and camera
42  // check for validity before proceeding
43  //----------------------------------------------------------------------------
44  uchar llid = (uchar)(get_local_id(0) + get_local_size(0)*get_local_id(1));
45  int i=0,j=0;
46  i=get_global_id(0);
47  j=get_global_id(1);
48
49  // check to see if the thread corresponds to an actual pixel as in some
50  // cases #of threads will be more than the pixels.
51  if (i>=(*exp_image_dims).z || j>=(*exp_image_dims).w)
52    return;
53
54  //Store image index (may save a register).  Also initialize VIS and expected_int
55  imIndex[llid] = j*get_global_size(0)+i;
56
57  //----------------------------------------------------------------------------
58  // Calculate ray origin, and direction
59  // (make sure ray direction is never axis aligned)
60  //----------------------------------------------------------------------------
61  float4 ray_o = ray_origins[ imIndex[llid] ];
62  float4 ray_d = ray_directions[ imIndex[llid] ];
63  float ray_ox, ray_oy, ray_oz, ray_dx, ray_dy, ray_dz;
64  //calc_scene_ray(linfo, camera, i, j, &ray_ox, &ray_oy, &ray_oz, &ray_dx, &ray_dy, &ray_dz);
65  calc_scene_ray_generic_cam(linfo, ray_o, ray_d, &ray_ox, &ray_oy, &ray_oz, &ray_dx, &ray_dy, &ray_dz);
66
67  //----------------------------------------------------------------------------
68  // we know i,j map to a point on the image, have calculated ray
69  // BEGIN RAY TRACE
70  //----------------------------------------------------------------------------
71
72  float4  expint = exp_image[imIndex[llid]];
73#ifdef YUV
74  expint = rgb2yuv(expint);
75#endif
76  float  vis     = vis_image[imIndex[llid]];
77  AuxArgs aux_args;
78  aux_args.alpha  = alpha_array;
79  aux_args.mog    = mixture_array;
80  aux_args.expint = &expint;
81  aux_args.vis    = &vis;
82  float nearplane = nearfarplanes[0]/linfo->block_len;
83  float farplane = nearfarplanes[1]/linfo->block_len;
84  cast_ray( i, j,
85            ray_ox, ray_oy, ray_oz,
86            ray_dx, ray_dy, ray_dz,
87            linfo, tree_array,                                    //scene info
88            local_tree, bit_lookup, cumsum, &vis, aux_args,nearplane,farplane);      //utility info
89
90  //store the expected intensity (as UINT)
91  //YUV edit
92#ifdef YUV
93  expint = yuv2rgb(expint);
94#endif
95  exp_image[imIndex[llid]] = expint;
96
97  //store visibility at the end of this block
98  vis_image[imIndex[llid]] = vis;
99}
100
101void step_cell_render(AuxArgs aux_args, int data_ptr, uchar llid, float d)
102{
103  float alpha = aux_args.alpha[data_ptr];
104  float diff_omega = exp(-alpha*d);
105  float4 expected_int_cell = 0.0f;
106
107  // for rendering only
108  if (diff_omega<0.995f)
109  {
110    uchar8 udata = as_uchar8(aux_args.mog[data_ptr]);
111    float8 data = convert_float8(udata)/255.0f;
112
113    //expected cell is just the means
114    expected_int_cell = (float4)data.s0123;
115
116    //undo the step taken in compress RGB U/V is in [0,1],
117    // put them back in ranges U in [-.436, .436] and V in [-.615, .615]
118#ifdef YUV
119    expected_int_cell.y = expected_int_cell.y*U_RANGE - U_MAX;
120    expected_int_cell.z = expected_int_cell.z*V_RANGE - V_MAX;
121#endif
122  }
123
124  //calc and store visibility
125  float omega = (*aux_args.vis) * (1.0f - diff_omega);
126  (*aux_args.vis) *= diff_omega;
127  (*aux_args.expint) += expected_int_cell*omega;
128}
129
130#endif // RENDER
131