Commit 86642e2b authored by karius's avatar karius

status quo push

parent fc8adb95
......@@ -92,8 +92,6 @@ void TransformationGrid::read_matrix_SO3_sampling(const char * file_name,float4
for (int i=0;i<c;i++){
h_quat_to_euler(h_rotations[i],curr_euler);
fho << curr_euler.x << "," << curr_euler.y << "," << curr_euler.z << "\n";
printf("%f,%f,%f,%f\n",h_rotations[i].x,h_rotations[i].y,h_rotations[i].z,h_rotations[i].w);
printf("%f,%f,%f\n",curr_euler.x,curr_euler.y,curr_euler.z);
}
fho.close();
}
......@@ -963,10 +961,29 @@ void float_min_max(float * d_data,size_t num_data, float * d_min, float * d_max)
if (blockIdx.x == 0 && threadIdx.x ==0){
d_min[0] = fminf(_min[0],_min[1]);
d_max[0] = fmaxf(_max[0],_max[1]);
printf("Min:%f %f\n",d_min[0],fminf(_min[0],_min[1]));
}
}
__host__
void TransformationGrid::shift_and_scale_chamfer_score(float * d_score, float * d_normalized_score,float *d_min, float *d_max, size_t num_scores){
// shift_and_scale_envelope_score_kernel<128><<<32,128>>>(d_score,d_normalized_score,d_min,d_max,num_scores);
float_min_max<128><<<32,128>>>(d_score,num_scores,d_min,d_max);
normalize_envelope unary_function;
unary_function.d_min = d_min;
unary_function.d_max = d_max;
thrust::transform(thrust::device_pointer_cast(&d_score[0]),thrust::device_pointer_cast(&d_score[num_scores-1]),thrust::device_pointer_cast(&d_normalized_score[0]),unary_function);
float * h_sum = (float *) malloc(sizeof(*h_sum));
// h_sum[0] = thrust::reduce(thrust::device_pointer_cast(&d_normalized_score[0]),thrust::device_pointer_cast(&d_normalized_score[num_scores-1]));
float sum = thrust::reduce(thrust::device_pointer_cast(&d_normalized_score[0]),thrust::device_pointer_cast(&d_normalized_score[num_scores-1]));
float * d_sum;
cudaMalloc((void **)&d_sum,sizeof(*d_sum));
cudaMemcpy(d_sum,h_sum,sizeof(*h_sum),cudaMemcpyHostToDevice);
CudaCheckError();
thrust::for_each(thrust::device_pointer_cast(&d_normalized_score[0]),thrust::device_pointer_cast(&d_normalized_score[num_scores-1]), thrust::placeholders::_1 /= sum);
thrust::for_each(thrust::device_pointer_cast(&d_normalized_score[0]),thrust::device_pointer_cast(&d_normalized_score[num_scores-1]), entropy());
float entr = -1.0*thrust::reduce(thrust::device_pointer_cast(&d_normalized_score[0]),thrust::device_pointer_cast(&d_normalized_score[num_scores-1]));
printf("Entropy for envelope score: %f\n",entr);
}
__host__
......
......@@ -56,6 +56,7 @@ public:
__host__ static void h_matrix_to_quat(float * const& m,float4 & quaternion);
__host__ static void read_matrix_SO3_sampling(const char * file_name,float4 *& h_rotations, uint *& h_num_rotations);
__host__ static void shift_and_scale_envelope_score(float * d_score, float * d_normalized_score,float *d_min, float *d_max, size_t num_scores);
__host__ static void shift_and_scale_chamfer_score(float * d_score, float * d_normalized_score,float *d_min, float *d_max, size_t num_scores);
TransformationGrid(const std::pair<float3,float3> & bounding_box_outer,const std::pair<float3,float3> & bounding_box_inner,
float * const& translation_offset,const float4 & angle_box, const float & alpha, const float & tau_in);
//Testgrid constructor
......
No preview for this file type
......@@ -315,7 +315,7 @@ void kai_score(float4 * d_coords, int * d_segmentations, uint * d_labels, float
while (b < grid.d_num_transformations[0]){
//transformation to shared
if (threadIdx.x == 0) grid.d_transformation_to_memory(b,&rotation,&translation,grid.d_translation_offset);
// if (threadIdx.x == 0) printf("%u: %f %f %f %f\n",b,rotation.x,rotation.y,rotation.z,rotation.w);
if (threadIdx.x == 0) printf("%u: %f %f %f %f:",b,rotation.x,rotation.y,rotation.z,rotation.w);
// if (threadIdx.x == 0) printf("%u: %f %f %f \n",b,translation.x,translation.y,translation.z);
__syncthreads();
//set segments to 0 - referring to padded volume
......@@ -415,6 +415,9 @@ void kai_score(float4 * d_coords, int * d_segmentations, uint * d_labels, float
if (threadIdx.x < 1) reduce[threadIdx.x] += reduce[threadIdx.x + 1]; __syncthreads();
grid.d_scores[b] = reduce[0];
b+=gridDim.x;
if (threadIdx.x == 0){
printf("%f:\n",reduce[0]);
}
}
}
......@@ -465,7 +468,7 @@ void kai_score(float4 * d_coords, int * d_segmentations, uint * d_labels, float
// uint n_rot_test = 30;
// std::pair<float3,float3> bounding_box_outer = density.get_bounding_box();
// std::pair<float3,float3> bounding_box_inner = pixel_coords.get_rotationally_safe_bounding_box();
// TransformationGrid transformation_grid(bounding_box_outer,bounding_box_inner,difference_map.h_mid_point,tau_in,n_rot_test);
// TransformationGrid transformation_grid(bounding_box_outer,bounding_box_inner,density.h_mid_point,tau_in,"rot_samples/N15_M1776_OctaC1.dat");
// printf("Transformation num: %u\n",transformation_grid.h_num_transformations[0]);
// printf("Translation num: %u\n",transformation_grid.h_num_translations[0]);
// printf("Translational shift: %f %f %f\n",transformation_grid.h_translation_step_shift->x,transformation_grid.h_translation_step_shift->y,transformation_grid.h_translation_step_shift->z);
......@@ -528,162 +531,164 @@ void kai_score(float4 * d_coords, int * d_segmentations, uint * d_labels, float
// nvmlShutdown();
// printf("\n\n");
//}
//
//BOOST_AUTO_TEST_CASE(envelope_test_simpler)
// {
// printf(" ============================= Testing Envelope score ==========================\n");
// printf(" ============================= ++++++++++++++++++++++ ==========================\n");
// //memory management
// nvmlInit_v2();
// //checkout free memory
// nvmlMemory_t memory_info;
// nvmlDevice_t device_handle;
// nvmlDeviceGetHandleByIndex_v2(0,&device_handle);
// nvmlDeviceGetMemoryInfo(device_handle, &memory_info);
//
// float above_value = -1.0;
// float below_value = 0.0;
// float coord_dim_outer[3] = {70,70,70};
// float coord_dim[3] = {50,50,50};
// float distance_threshold = 0.5;
// uint pixel_dim_outer[3] = {70,70,70};
// uint pixel_dim[3] = {50,50,50};
// uint3 cube_pixel_dim= {30,30,30};
// Density density;
// Density::from_bounds(density, &coord_dim_outer[0],&pixel_dim_outer[0]);
//// Density density = Density::from_mrc("test/tauA_Brf1_TBP_negstain_30A.mrc");
// density.init_device();
// density.to_device();
// density.mid_cube_with_surface(cube_pixel_dim,1,1.0,1.5);
//
// Density binarized_map;
// density.cut_and_binarize(binarized_map,1.0,above_value,below_value,{20,20,20});
//
// Density to_fit_density;
// Density::from_bounds(to_fit_density, &coord_dim[0],&pixel_dim[0]);
// to_fit_density.mid_cube_with_surface(cube_pixel_dim,1,1.0,1.5);
//
// //every pixel a particle, for test reasons
// thrust::device_vector<uint> td_on_surface_to_fit;
// size_t on_surface_num_to_fit;
// to_fit_density.indices_in_range(td_on_surface_to_fit,true,1.25,2,on_surface_num_to_fit);
//
// Particles pixel_coords = to_fit_density.particles_from_density(td_on_surface_to_fit);
// pixel_coords.origin_to_center();
//
// float memory_usage_factor = 0.5;
// size_t density_memory_usage = pixel_coords.particle_count()*sizeof(uint) + binarized_map.h_pixel_vol()*sizeof(float);
// unsigned long int max = 100;
// uint work_spaces = std::min(max,((uint)(((float)memory_info.free)*memory_usage_factor))/density_memory_usage);
//
// //set up TransformationGrid
// float tau_in = 1.0;
// std::pair<float3,float3> bounding_box_outer = density.get_bounding_box();
// std::pair<float3,float3> bounding_box_inner = pixel_coords.get_rotationally_safe_bounding_box();
// TransformationGrid transformation_grid(bounding_box_outer,bounding_box_inner,binarized_map.h_mid_point,tau_in,"rot_samples/N15_M1776_OctaC1.dat");
// CudaCheckError();
// printf("Transformation num: %u\n",transformation_grid.h_num_transformations[0]);
// printf("Translation num: %u\n",transformation_grid.h_num_translations[0]);
// printf("Translational shift: %f %f %f\n",transformation_grid.h_translation_step_shift->x,transformation_grid.h_translation_step_shift->y,transformation_grid.h_translation_step_shift->z);
// printf("Rotation num: %u\n",transformation_grid.h_num_rotations[0]);
//
//// TransformationGrid transformation_grid(binarized_map.h_mid_point);
// const size_t envelope_score_block_dim = 256;
// size_t envelope_score_grid_dim = work_spaces;
// size_t envelope_workspace_replaced_indices_volume = work_spaces*pixel_coords.particle_count()*sizeof(uint);
// size_t envelope_workspace_replaced_values_buffer_volume = work_spaces*pixel_coords.particle_count()*sizeof(float);
// //allocate workspace resources
//
// float * d_replaced_value_buffer;
// cudaMalloc((void **)&d_replaced_value_buffer,envelope_workspace_replaced_values_buffer_volume);
// uint * d_replaced_linear_indices;
// cudaMalloc((void **)&d_replaced_linear_indices,envelope_workspace_replaced_indices_volume);
// CudaCheckError();
// float * d_workspaces;
// cudaMalloc((void **)&d_workspaces,work_spaces*binarized_map.h_pixel_vol()*sizeof(float));
// printf("Allocating %u workspaces for %u pixels each occupying %u bytes of memory ...\n",work_spaces,binarized_map.h_pixel_vol(),envelope_workspace_replaced_indices_volume + envelope_workspace_replaced_values_buffer_volume);
// CudaCheckError();
// //(Density binarized_map, float * d_workspaces, uint * d_linear_offsets, float4 * d_coords,size_t num_particles , TransformationGrid grid)
// envelope_score<envelope_score_block_dim><<<envelope_score_grid_dim,envelope_score_block_dim>>>
// (binarized_map, d_workspaces,d_replaced_value_buffer, d_replaced_linear_indices,pixel_coords.d_data,pixel_coords.particle_count(),transformation_grid);
//// transformation_grid.write_to_csv("test/env_score.csv");
// float * d_normalized_scores;
// cudaMalloc((void **)&d_normalized_scores,transformation_grid.h_num_transformations[0]*sizeof(*d_normalized_scores));
// float * d_min;
// cudaMalloc((void **)&d_min,sizeof(*d_min));
// float * d_max;
// cudaMalloc((void **)&d_max,sizeof(*d_max));
// transformation_grid.shift_and_scale_envelope_score(transformation_grid.d_scores,d_normalized_scores,d_min,d_max,transformation_grid.h_num_transformations[0]);
// transformation_grid.write_to_csv("test/env_score.csv");
// printf("Finished testing envelope score ... \n\n");
//}
BOOST_AUTO_TEST_CASE(envelope_test_simpler)
{
printf(" ============================= Testing Envelope score ==========================\n");
printf(" ============================= ++++++++++++++++++++++ ==========================\n");
//memory management
nvmlInit_v2();
//checkout free memory
nvmlMemory_t memory_info;
nvmlDevice_t device_handle;
nvmlDeviceGetHandleByIndex_v2(0,&device_handle);
nvmlDeviceGetMemoryInfo(device_handle, &memory_info);
float above_value = -1.0;
float below_value = 0.0;
float coord_dim_outer[3] = {70,70,70};
float coord_dim[3] = {50,50,50};
float distance_threshold = 0.5;
uint pixel_dim_outer[3] = {70,70,70};
uint pixel_dim[3] = {50,50,50};
uint3 cube_pixel_dim= {30,30,30};
Density density;
Density::from_bounds(density, &coord_dim_outer[0],&pixel_dim_outer[0]);
// Density density = Density::from_mrc("test/tauA_Brf1_TBP_negstain_30A.mrc");
density.init_device();
density.to_device();
density.mid_cube_with_surface(cube_pixel_dim,1,1.0,1.5);
BOOST_AUTO_TEST_CASE(chamfer_test_simpler)
{
printf(" ============================= Testing Chamfer score ===========================\n");
printf(" ============================= +++++++++++++++++++++ ===========================\n");
Density binarized_map;
density.cut_and_binarize(binarized_map,1.0,above_value,below_value,{20,20,20});
//memory management
nvmlInit_v2();
//checkout free memory
nvmlMemory_t memory_info;
nvmlDevice_t device_handle;
nvmlDeviceGetHandleByIndex_v2(0,&device_handle);
nvmlDeviceGetMemoryInfo(device_handle, &memory_info);
Density to_fit_density;
Density::from_bounds(to_fit_density, &coord_dim[0],&pixel_dim[0]);
to_fit_density.mid_cube_with_surface(cube_pixel_dim,1,1.0,1.5);
float coord_dim_outer[3] = {70,70,70};
float coord_dim[3] = {50,50,50};
float distance_threshold = 0.5;
uint pixel_dim_outer[3] = {70,70,70};
uint pixel_dim[3] = {50,50,50};
uint3 cube_pixel_dim= {30,30,30};
Density density;
Density::from_bounds(density,&coord_dim_outer[0],&pixel_dim_outer[0]);
density.init_device();
density.to_device();
density.mid_cube_with_surface(cube_pixel_dim,1,1.0,1.5);
//every pixel a particle, for test reasons
thrust::device_vector<uint> td_on_surface_to_fit;
size_t on_surface_num_to_fit;
to_fit_density.indices_in_range(td_on_surface_to_fit,true,1.25,2,on_surface_num_to_fit);
Density to_fit_density;
Density::from_bounds(to_fit_density,&coord_dim[0],&pixel_dim[0]);
to_fit_density.mid_cube_with_surface(cube_pixel_dim,1,1.0,1.5);
Particles pixel_coords = to_fit_density.particles_from_density(td_on_surface_to_fit);
pixel_coords.origin_to_center();
thrust::device_vector<uint> td_on_surface;
size_t on_surface_num;
density.indices_in_range(td_on_surface,true,1.25,2,on_surface_num);
float memory_usage_factor = 0.5;
size_t density_memory_usage = pixel_coords.particle_count()*sizeof(uint) + binarized_map.h_pixel_vol()*sizeof(float);
unsigned long int max = 100;
uint work_spaces = std::min(max,((uint)(((float)memory_info.free)*memory_usage_factor))/density_memory_usage);
thrust::device_vector<uint> td_on_surface_to_fit;
size_t on_surface_num_to_fit;
to_fit_density.indices_in_range(td_on_surface_to_fit,true,1.25,2,on_surface_num_to_fit);
//set up TransformationGrid
float tau_in = 1.0;
std::pair<float3,float3> bounding_box_outer = density.get_bounding_box();
std::pair<float3,float3> bounding_box_inner = pixel_coords.get_rotationally_safe_bounding_box();
TransformationGrid transformation_grid(bounding_box_outer,bounding_box_inner,binarized_map.h_mid_point,tau_in,"rot_samples/N23_M5880_IcoC7.dat");
CudaCheckError();
printf("Transformation num: %u\n",transformation_grid.h_num_transformations[0]);
printf("Translation num: %u\n",transformation_grid.h_num_translations[0]);
printf("Translational shift: %f %f %f\n",transformation_grid.h_translation_step_shift->x,transformation_grid.h_translation_step_shift->y,transformation_grid.h_translation_step_shift->z);
printf("Rotation num: %u\n",transformation_grid.h_num_rotations[0]);
Particles pixel_coords = to_fit_density.particles_from_density(td_on_surface_to_fit);
pixel_coords.origin_to_center();
// TransformationGrid transformation_grid(binarized_map.h_mid_point);
const size_t envelope_score_block_dim = 256;
size_t envelope_score_grid_dim = work_spaces;
size_t envelope_workspace_replaced_indices_volume = work_spaces*pixel_coords.particle_count()*sizeof(uint);
size_t envelope_workspace_replaced_values_buffer_volume = work_spaces*pixel_coords.particle_count()*sizeof(float);
//allocate workspace resources
Density difference_map;
density.difference_map_density_threshold(difference_map,1.4,1.6,td_on_surface);
CudaCheckError();
float * d_replaced_value_buffer;
cudaMalloc((void **)&d_replaced_value_buffer,envelope_workspace_replaced_values_buffer_volume);
uint * d_replaced_linear_indices;
cudaMalloc((void **)&d_replaced_linear_indices,envelope_workspace_replaced_indices_volume);
CudaCheckError();
float * d_workspaces;
cudaMalloc((void **)&d_workspaces,work_spaces*binarized_map.h_pixel_vol()*sizeof(float));
printf("Allocating %u workspaces for %u pixels each occupying %u bytes of memory ...\n",work_spaces,binarized_map.h_pixel_vol(),envelope_workspace_replaced_indices_volume + envelope_workspace_replaced_values_buffer_volume);
CudaCheckError();
//(Density binarized_map, float * d_workspaces, uint * d_linear_offsets, float4 * d_coords,size_t num_particles , TransformationGrid grid)
envelope_score<envelope_score_block_dim><<<envelope_score_grid_dim,envelope_score_block_dim>>>
(binarized_map, d_workspaces,d_replaced_value_buffer, d_replaced_linear_indices,pixel_coords.d_data,pixel_coords.particle_count(),transformation_grid);
// transformation_grid.write_to_csv("test/env_score.csv");
float * d_normalized_scores;
cudaMalloc((void **)&d_normalized_scores,transformation_grid.h_num_transformations[0]*sizeof(*d_normalized_scores));
float * d_min;
cudaMalloc((void **)&d_min,sizeof(*d_min));
float * d_max;
cudaMalloc((void **)&d_max,sizeof(*d_max));
transformation_grid.shift_and_scale_envelope_score(transformation_grid.d_scores,d_normalized_scores,d_min,d_max,transformation_grid.h_num_transformations[0]);
printf("Finished testing envelope score ... \n\n");
}
float tau_in = 1.0;
uint n_rot_test = 30;
std::pair<float3,float3> bounding_box_outer = density.get_bounding_box();
std::pair<float3,float3> bounding_box_inner = pixel_coords.get_rotationally_safe_bounding_box();
TransformationGrid transformation_grid(bounding_box_outer,bounding_box_inner,density.h_mid_point,tau_in,"rot_samples/N15_M1776_OctaC1.dat");
printf("Transformation num: %u\n",transformation_grid.h_num_transformations[0]);
printf("Translation num: %u\n",transformation_grid.h_num_translations[0]);
printf("Translational shift: %f %f %f\n",transformation_grid.h_translation_step_shift->x,transformation_grid.h_translation_step_shift->y,transformation_grid.h_translation_step_shift->z);
printf("Rotation num: %u\n",transformation_grid.h_num_rotations[0]);
//BOOST_AUTO_TEST_CASE(chamfer_test_simpler)
//{
// printf(" ============================= Testing Chamfer score ===========================\n");
// printf(" ============================= +++++++++++++++++++++ ===========================\n");
//
// float coord_dim_outer[3] = {70,70,70};
// float coord_dim[3] = {50,50,50};
// float distance_threshold = 0.5;
// uint pixel_multiplier = 3;
// uint pixel_dim_outer[3] = {70,70,70};
// uint pixel_dim[3] = {50,50,50};
// uint3 cube_pixel_dim= {30,30,30};
// for (int i=0; i<3; i++)pixel_dim_outer[i] *= pixel_multiplier;
// for (int i=0; i<3; i++)pixel_dim[i] *= pixel_multiplier;
// cube_pixel_dim.x *= pixel_multiplier;
// cube_pixel_dim.y *= pixel_multiplier;
// cube_pixel_dim.z *= pixel_multiplier;
//
// Density density;
// Density::from_bounds(density,&coord_dim_outer[0],&pixel_dim_outer[0]);
// density.init_device();
// density.to_device();
// density.mid_cube_with_surface(cube_pixel_dim,1,1.0,1.5);
//
// Density to_fit_density;
// Density::from_bounds(to_fit_density,&coord_dim[0],&pixel_dim[0]);
// to_fit_density.mid_cube_with_surface(cube_pixel_dim,1,1.0,1.5);
//
// thrust::device_vector<uint> td_on_surface;
// size_t on_surface_num;
// density.indices_in_range(td_on_surface,true,1.25,2,on_surface_num);
//
// thrust::device_vector<uint> td_on_surface_to_fit;
// size_t on_surface_num_to_fit;
// to_fit_density.indices_in_range(td_on_surface_to_fit,true,1.25,2,on_surface_num_to_fit);
//
// Particles pixel_coords = to_fit_density.particles_from_density(td_on_surface_to_fit);
// pixel_coords.origin_to_center();
//
// Density difference_map;
// density.difference_map_density_threshold(difference_map,1.4,1.6,td_on_surface);
// CudaCheckError();
//
// float tau_in = 1.0;
// uint n_rot_test = 30;
// std::pair<float3,float3> bounding_box_outer = density.get_bounding_box();
// std::pair<float3,float3> bounding_box_inner = pixel_coords.get_rotationally_safe_bounding_box();
// TransformationGrid transformation_grid(bounding_box_outer,bounding_box_inner,difference_map.h_mid_point,tau_in,n_rot_test);
// printf("Transformation num: %u\n",transformation_grid.h_num_transformations[0]);
// printf("Translation num: %u\n",transformation_grid.h_num_translations[0]);
// printf("Translational shift: %f %f %f\n",transformation_grid.h_translation_step_shift->x,transformation_grid.h_translation_step_shift->y,transformation_grid.h_translation_step_shift->z);
// printf("Rotation num: %u\n",transformation_grid.h_num_rotations[0]);
//
// size_t chamfer_kernel_block_dim = 128;
// size_t chamfer_kernel_grid_dim = 32;
// size_t chamfer_kernel_shared_mem = sizeof(float4) + sizeof(float4);
// auto start = std::chrono::high_resolution_clock::now();
// chamfer_score<<<chamfer_kernel_grid_dim,chamfer_kernel_block_dim,chamfer_kernel_shared_mem>>>
// (pixel_coords.d_data,difference_map,pixel_coords.particle_count(),transformation_grid);
// cudaDeviceSynchronize();
// auto finish = std::chrono::high_resolution_clock::now();
// std::chrono::duration<double> elapsed = finish - start;
// std::cout << "Elapsed time: " << elapsed.count() << " s\n";
size_t chamfer_kernel_block_dim = 128;
size_t chamfer_kernel_grid_dim = 32;
size_t chamfer_kernel_shared_mem = sizeof(float4) + sizeof(float4);
auto start = std::chrono::high_resolution_clock::now();
chamfer_score<<<chamfer_kernel_grid_dim,chamfer_kernel_block_dim,chamfer_kernel_shared_mem>>>
(pixel_coords.d_data,difference_map,pixel_coords.particle_count(),transformation_grid);
cudaDeviceSynchronize();
auto finish = std::chrono::high_resolution_clock::now();
std::chrono::duration<double> elapsed = finish - start;
std::cout << "Elapsed time: " << elapsed.count() << " s\n";
// transformation_grid.normalize_chamfer_score();
// transformation_grid.write_to_csv("test/chamfer_score.csv");
// printf("\n\n");
//}
transformation_grid.write_to_csv("test/chamfer_score.csv");
printf("\n\n");
}
//BOOST_AUTO_TEST_CASE(ccl_test)
......
This source diff could not be displayed because it is too large. You can view the blob instead.
No preview for this file type
No preview for this file type
Markdown is supported
0% or .
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment