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