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