37 #include <boost/filesystem.hpp> 38 #include <boost/shared_array.hpp> 43 #define CL_HPP_TARGET_OPENCL_VERSION 120 44 #define CL_USE_DEPRECATED_OPENCL_1_2_APIS 45 #define __CL_ENABLE_EXCEPTIONS 47 #if defined(__APPLE__) || defined(__MACOSX) 48 #include <OpenCL/opencl.h> 54 #define MAX_SOURCE_SIZE (0x1024) 59 typedef boost::shared_array<float>
floatArr;
61 using Vec = BaseVector<float>;
66 ClSOR(floatArr& points,
size_t num_points,
int k,
int device = 0);
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" 145 " unsigned int pos = 0; \n" 146 " unsigned int current_dim = 0; \n" 147 " while(pos < num_splits) \n" 149 " current_dim = (unsigned int)(D_kd_tree_splits[pos]); \n" 150 " if(current_dim == 0) \n" 152 " if(x <= D_kd_tree_values[pos] ) \n" 158 " } else if(current_dim == 1) { \n" 159 " if(y <= D_kd_tree_values[pos] ){ \n" 165 " if(z <= D_kd_tree_values[pos] ){ \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" 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" 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" 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" 204 " correct = num_splits - start; \n" 205 " }else if(end > num_values) \n" 207 " correct = num_values - end; \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" 217 " nearest_index = (unsigned int)(D_kd_tree_values[i]+ 0.5); \n" 218 " if(nearest_index < num_points) \n" 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" 231 " D_distances[tid] = mean/k;"
Datastructures for holding loaded data.
ClSOR(floatArr &points, size_t num_points, int k, int device=0)
QueryPoint< Vec > QueryPointC
cl_uint m_threads_per_block
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
boost::shared_array< float > floatArr
int getInliers(lvr2::indexArray &inliers)
cl_ulong m_device_global_memory
const char * SOR_KERNEL_STRING
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
LBPointArray< float > * kd_tree_values