1// Ingest label map kernel, 2// MOG type: boxm2_label_short 3// assign the label of the ray to all voxels along the ray, when rendering boxm2_label_short, low prob ones aren't displayed anyways 4 5#ifdef INGEST_LABEL_MAP 6 7//need to define a struct of type AuxArgs with auxiliary arguments 8// to supplement cast ray args 9typedef struct 10{ 11 __global MOG_TYPE * label_data; 12 __global uchar label; 13} AuxArgs; 14 15/* previous when label is directly put to the voxel 16void step_cell_ingest_label_map(AuxArgs aux_args, int data_ptr) 17{ 18 aux_args.label_data[data_ptr] = (MOG_TYPE)aux_args.label; 19}*/ 20 21 22/* NEW ingest label step cell, keep the previous value in the upper two bits (bit 7 and bit 6) */ 23void step_cell_ingest_label_map(AuxArgs aux_args, int data_ptr) 24{ 25 MOG_TYPE current = aux_args.label_data[data_ptr]; 26 MOG_TYPE invalid = 100; 27 if (current == invalid) // invalid 28 current = 0; 29 else if (current >= 2) // 0 is invalid, 1 is horizontal, 2 and up are vertical 30 current = 2; // make all types of vertical 2, so we have 3 values as orientation 31 //current << 6; // shift 6 bits to the left, so upper 2 bits are orientation bits 32 current *= 64; 33 current += (MOG_TYPE)aux_args.label; // add the land type, so the lower 6 bits will represent the land type 34 35 aux_args.label_data[data_ptr] = current; 36} 37 38//forward declare cast ray (so you can use it) 39void cast_ray(int,int,float,float,float,float,float,float, 40 __constant RenderSceneInfo*, __global int4*, 41 __local uchar16*, __constant uchar *,__local uchar *, 42 float*, AuxArgs,float tnear, float tfar); 43__kernel 44void ingest_label_map(__constant RenderSceneInfo * linfo, 45 __global uint4 * image_dims, 46 __global float4 * ray_origin_buff, 47 __global int4 * tree_array, 48 __global ushort * data_array, 49 __global uchar * data_buff, 50 __constant uchar * bit_lookup, 51 __local uchar16 * local_tree, 52 __local uchar * cumsum, // cumulative sum helper for data pointer 53 __local int * imIndex) 54{ 55 //---------------------------------------------------------------------------- 56 //get local id (0-63 for an 8x8) of this patch + image coordinates and camera 57 // check for validity before proceeding 58 //---------------------------------------------------------------------------- 59 uchar llid = (uchar)(get_local_id(0) + get_local_size(0)*get_local_id(1)); 60 int i=0,j=0; 61 i=get_global_id(0); 62 j=get_global_id(1); 63 imIndex[llid] = j*get_global_size(0)+i; 64 // check to see if the thread corresponds to an actual pixel as in some 65 // cases #of threads will be more than the pixels. 66 if (i>=(*image_dims).z || j>=(*image_dims).w) 67 return; 68 69 //---------------------------------------------------------------------------- 70 // Calculate ray origin, and direction 71 // (make sure ray direction is never axis aligned) 72 //---------------------------------------------------------------------------- 73 float4 ray_o = ray_origin_buff[ imIndex[llid] ]; 74 uchar data = data_buff[ imIndex[llid] ]; 75 76 float4 ray_d = (float4)( 0.001, 0.001, -1.0, 1.0); 77 78 float ray_ox = 0.0f;float ray_oy = 0.0f;float ray_oz = 0.0f; 79 float ray_dx = 0.0f;float ray_dy = 0.0f;float ray_dz = 0.0f; 80 81 calc_scene_ray_generic_cam(linfo, ray_o, ray_d, 82 &ray_ox, &ray_oy, &ray_oz, 83 &ray_dx, &ray_dy, &ray_dz); 84 85 ////---------------------------------------------------------------------------- 86 //// we know i,j map to a point on the image, have calculated ray 87 //// BEGIN RAY TRACE 88 ////---------------------------------------------------------------------------- 89 90 AuxArgs aux_args; 91 //float out1 =0.0f; 92 aux_args.label_data = data_array; 93 //aux_args.outimg = out1; 94 aux_args.label = data; 95 float vis =1.0; 96 97 cast_ray( i, j, 98 ray_ox, ray_oy, ray_oz, 99 ray_dx, ray_dy, ray_dz, 100 linfo, tree_array, //scene info 101 local_tree, bit_lookup, cumsum, &vis, aux_args,0,MAXFLOAT); //utility info 102 103} 104 105 106__kernel 107void ingest_label_map_with_dir(__constant RenderSceneInfo * linfo, 108 __global uint4 * image_dims, 109 __global float4 * ray_origin_buff, 110 __global float4 * ray_dir_buff, 111 __global int4 * tree_array, 112 __global ushort * data_array, 113 __global uchar * data_buff, 114 __constant uchar * bit_lookup, 115 __local uchar16 * local_tree, 116 __local uchar * cumsum, // cumulative sum helper for data pointer 117 __local int * imIndex) 118{ 119 //---------------------------------------------------------------------------- 120 //get local id (0-63 for an 8x8) of this patch + image coordinates and camera 121 // check for validity before proceeding 122 //---------------------------------------------------------------------------- 123 uchar llid = (uchar)(get_local_id(0) + get_local_size(0)*get_local_id(1)); 124 int i=0,j=0; 125 i=get_global_id(0); 126 j=get_global_id(1); 127 imIndex[llid] = j*get_global_size(0)+i; 128 // check to see if the thread corresponds to an actual pixel as in some 129 // cases #of threads will be more than the pixels. 130 if (i>=(*image_dims).z || j>=(*image_dims).w) 131 return; 132 133 //---------------------------------------------------------------------------- 134 // Calculate ray origin, and direction 135 // (make sure ray direction is never axis aligned) 136 //---------------------------------------------------------------------------- 137 float4 ray_o = ray_origin_buff[ imIndex[llid] ]; 138 uchar data = data_buff[ imIndex[llid] ]; 139 140 float4 ray_d = ray_dir_buff[ imIndex[llid] ]; 141 142 float ray_ox = 0.0f;float ray_oy = 0.0f;float ray_oz = 0.0f; 143 float ray_dx = 0.0f;float ray_dy = 0.0f;float ray_dz = 0.0f; 144 145 calc_scene_ray_generic_cam(linfo, ray_o, ray_d, 146 &ray_ox, &ray_oy, &ray_oz, 147 &ray_dx, &ray_dy, &ray_dz); 148 149 ////---------------------------------------------------------------------------- 150 //// we know i,j map to a point on the image, have calculated ray 151 //// BEGIN RAY TRACE 152 ////---------------------------------------------------------------------------- 153 154 AuxArgs aux_args; 155 //float out1 =0.0f; 156 aux_args.label_data = data_array; 157 //aux_args.outimg = out1; 158 aux_args.label = data; 159 float vis =1.0; 160 161 cast_ray( i, j, 162 ray_ox, ray_oy, ray_oz, 163 ray_dx, ray_dy, ray_dz, 164 linfo, tree_array, //scene info 165 local_tree, bit_lookup, cumsum, &vis, aux_args,0,MAXFLOAT); //utility info 166 167} 168 169 170#endif // INGEST_LABEL_MAP 171