ClStatisticalOutlierFilter.hpp
Go to the documentation of this file.
1 
28 #ifndef __CLSOR_H
29 #define __CLSOR_H
30 
35 #include "lvr2/io/DataStruct.hpp"
36 
37 #include <boost/filesystem.hpp>
38 #include <boost/shared_array.hpp>
39 #include <string.h>
40 #include <stdio.h>
41 #include <stdlib.h>
42 
43 #define CL_HPP_TARGET_OPENCL_VERSION 120
44 #define CL_USE_DEPRECATED_OPENCL_1_2_APIS
45 #define __CL_ENABLE_EXCEPTIONS
46 
47 #if defined(__APPLE__) || defined(__MACOSX)
48  #include <OpenCL/opencl.h>
49 #else
50  #include <CL/cl.h>
51 #endif
53 
54 #define MAX_SOURCE_SIZE (0x1024)
55 
56 namespace lvr2
57 {
58 
59 typedef boost::shared_array<float> floatArr;
60 
61 using Vec = BaseVector<float>;
62 typedef QueryPoint<Vec> QueryPointC;
63 
64 class ClSOR {
65 public:
66  ClSOR(floatArr& points, size_t num_points, int k, int device = 0);
67  ~ClSOR();
68 
69 
77  void setK(int k);
78  void setMult(float std_dev_mult) { m_mult_ = std_dev_mult;};
79 
80  void calcDistances();
81  void calcStatistics();
82 
83  int getInliers(lvr2::indexArray& inliers);
84 
85  void freeGPU();
86 
87 private:
88 
89  void init();
90 
91  const char *getErrorString(cl_int error);
92 
93  void initKdTree();
94 
95  void getDeviceInformation(int platform_id=0, int device_id=0);
96 
97  void loadSORKernel();
98 
99  void initCl();
100 
101  void finalizeCl();
102 
103  //
104  int m_k;
105  double m_mult_;
106  double m_mean_;
107  double m_std_dev_;
108  // V->points and normals
112 
113 // LBPointArray<float> Result_Normals;
115  boost::shared_ptr<LBKdTree> kd_tree_gen;
116 
117 
118 
119 
120  // Device Information
121  cl_platform_id m_platform_id;
122  cl_device_id m_device_id;
123  cl_uint m_mps;
126  cl_int m_ret;
127  cl_context m_context;
128  cl_command_queue m_command_queue;
129  cl_program m_program_es;
130  cl_program m_program_in;
131  cl_kernel m_kernel_sor;
132 // cl_kernel m_kernel_normal_interpolation;
133 
134  cl_mem D_V;
137  cl_mem D_Distances;
138 
139 
140 const char *SOR_KERNEL_STRING = "\n"
141 "unsigned int GetKdTreePosition(__global const float* D_kd_tree_values,"
142 "const unsigned int num_values,__global const unsigned char* D_kd_tree_splits,"
143 "const unsigned int num_splits, float x, float y, float z) \n"
144 "{ \n"
145 " unsigned int pos = 0; \n"
146 " unsigned int current_dim = 0; \n"
147 " while(pos < num_splits) \n"
148 " { \n"
149 " current_dim = (unsigned int)(D_kd_tree_splits[pos]); \n"
150 " if(current_dim == 0) \n"
151 " { \n"
152 " if(x <= D_kd_tree_values[pos] ) \n"
153 " { \n"
154 " pos = pos*2+1; \n"
155 " } else { \n"
156 " pos = pos*2+2; \n"
157 " } \n"
158 " } else if(current_dim == 1) { \n"
159 " if(y <= D_kd_tree_values[pos] ){ \n"
160 " pos = pos*2+1; \n"
161 " }else{ \n"
162 " pos = pos*2+2; \n"
163 " } \n"
164 " } else { \n"
165 " if(z <= D_kd_tree_values[pos] ){ \n"
166 " pos = pos*2+1; \n"
167 " }else{ \n"
168 " pos = pos*2+2; \n"
169 " } \n"
170 " } \n"
171 " } \n"
172 " return pos; \n"
173 "} \n"
174 " \n"
175 "__kernel void SORKernel(__global const float* D_V, const unsigned unsigned int num_points,"
176 "__global const float* D_kd_tree_values, const unsigned int num_values,"
177 "__global const unsigned char* D_kd_tree_splits , const unsigned int num_splits,"
178 "__global float* D_distances, const unsigned int k) \n"
179 "{ \n"
180 " unsigned int loc_id = get_local_id(0); \n"
181 " unsigned int loc_size = get_local_size(0); \n"
182 " unsigned int glob_id = get_global_id(0); \n"
183 " unsigned int glob_size = get_global_size(0); \n"
184 " unsigned int group_id = get_group_id(0); \n"
185 " unsigned int group_size = get_num_groups(0); \n"
186 " unsigned int tid = glob_id; \n"
187 " const unsigned int offset = glob_size; \n"
188 " for(;tid < num_points; tid += offset) \n"
189 " { \n"
190 " unsigned int pos = GetKdTreePosition(D_kd_tree_values, "
191 "num_values,D_kd_tree_splits, num_splits,D_V[tid * 3], D_V[tid * 3 + 1], D_V[tid * 3 +2] ); \n"
192 " unsigned int vertex_index = (unsigned int)(D_kd_tree_values[pos]+ 0.5); \n"
193 " if(vertex_index < num_points) \n"
194 " { \n"
195 " float vertex_x = D_V[ vertex_index * 3 + 0 ]; \n"
196 " float vertex_y = D_V[ vertex_index * 3 + 1 ]; \n"
197 " float vertex_z = D_V[ vertex_index * 3 + 2 ]; \n"
198 " unsigned int nearest_index; \n"
199 " int start = pos-(k/2); \n"
200 " int end = pos+((k+1)/2); \n"
201 " int correct = 0; \n"
202 " if(start < num_splits) \n"
203 " { \n"
204 " correct = num_splits - start; \n"
205 " }else if(end > num_values) \n"
206 " { \n"
207 " correct = num_values - end; \n"
208 " } \n"
209 " start += correct; \n"
210 " end += correct; \n"
211 " float mean = 0.0; \n"
212 " unsigned int j = 0; \n"
213 " for(unsigned int i = start; i < end && i<num_values; i++ ) \n"
214 " { \n"
215 " if(i != pos) \n"
216 " { \n"
217 " nearest_index = (unsigned int)(D_kd_tree_values[i]+ 0.5); \n"
218 " if(nearest_index < num_points) \n"
219 " { \n"
220 " "
221 " // calculate distances and mean \n"
222 " float diff_x = D_V[ nearest_index * 3 + 0 ] - vertex_x; \n"
223 " float diff_y = D_V[ nearest_index * 3 + 1 ] - vertex_y; \n"
224 " float diff_z = D_V[ nearest_index * 3 + 2 ] - vertex_z; \n"
225 " float dist = sqrt((pow(diff_x, 2) + pow(diff_y, 2) + pow(diff_z, 2)));\n"
226 " mean += dist; \n"
227 " j++;\n"
228 " } \n"
229 " } \n"
230 " } \n"
231 " D_distances[tid] = mean/k;"
232 " } \n"
233 " } \n"
234 "} \n";
235 
236 
237 };
238 
239 } /* namespace lvr2 */
240 
241 #endif // !__ClSOR_H
Datastructures for holding loaded data.
ClSOR(floatArr &points, size_t num_points, int k, int device=0)
LBPointArray< float > V
QueryPoint< Vec > QueryPointC
Definition: CudaSurface.hpp:75
const char * getErrorString(cl_int error)
LBPointArray< unsigned char > * kd_tree_splits
void setMult(float std_dev_mult)
void initKdTree()
PRIVATE ///.
cl_command_queue m_command_queue
void setK(int k)
Set the number of k nearest neighbors k-neighborhood.
cl_platform_id m_platform_id
void getDeviceInformation(int platform_id=0, int device_id=0)
boost::shared_array< unsigned int > indexArray
Definition: DataStruct.hpp:128
boost::shared_array< float > floatArr
Definition: DataStruct.hpp:133
int getInliers(lvr2::indexArray &inliers)
KF_EXPORTS void error(const char *error_string, const char *file, const int line, const char *func="")
Error handler. All GPU functions from this subsystem call the function to report an error...
boost::shared_ptr< LBKdTree > kd_tree_gen
LBPointArray< float > m_distances
BaseVector< float > Vec
LBPointArray< float > * kd_tree_values


lvr2
Author(s): Thomas Wiemann , Sebastian Pütz , Alexander Mock , Lars Kiesow , Lukas Kalbertodt , Tristan Igelbrink , Johan M. von Behren , Dominik Feldschnieders , Alexander Löhr
autogenerated on Mon Feb 28 2022 22:46:06