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