Commit d83ae270 authored by karius's avatar karius

fixed sampling

parent 3728892c
......@@ -389,6 +389,7 @@ Density::Density(){
h_upper_right = (float *) malloc(3*sizeof(float));
h_mid_point = (float *) malloc(3*sizeof(float));
h_pixel_size = (float *) malloc(sizeof(float));
h_padding_dim = (uint *) malloc(3*sizeof(uint));
cudaMalloc((void **) &d_coord_dim,3*sizeof(float));
cudaMalloc((void **) &d_pixel_dim,3*sizeof(uint));
......@@ -396,6 +397,7 @@ Density::Density(){
cudaMalloc((void **) &d_upper_right,3*sizeof(float));
cudaMalloc((void **) &d_mid_point,3*sizeof(float));
cudaMalloc((void **) &d_pixel_size,sizeof(float));
cudaMalloc((void **) &d_padding_dim,3*sizeof(uint));
}
......@@ -428,6 +430,7 @@ Density::Density(float * const & coord_dim, uint * const & stamp_dim, const floa
th_data.resize(h_padding_dim[0]*h_padding_dim[1]*h_padding_dim[2]);
h_data = thrust::raw_pointer_cast(&th_data[0]);
host_initiated = true;
CudaCheckError();
}
Density::Density(float * const & coord_dim, const float & pixel_size){
......@@ -851,10 +854,15 @@ void Density::to_device(){
d_data = thrust::raw_pointer_cast(&td_data[0]);
}
cudaMemcpy(d_coord_dim,h_coord_dim,3*sizeof(float),cudaMemcpyHostToDevice);
CudaCheckError();
cudaMemcpy(d_pixel_dim,h_pixel_dim,3*sizeof(uint),cudaMemcpyHostToDevice);
CudaCheckError();
cudaMemcpy(d_lower_left,h_lower_left,3*sizeof(float),cudaMemcpyHostToDevice);
CudaCheckError();
cudaMemcpy(d_upper_right,h_upper_right,3*sizeof(float),cudaMemcpyHostToDevice);
CudaCheckError();
cudaMemcpy(d_mid_point,h_mid_point,3*sizeof(float),cudaMemcpyHostToDevice);
CudaCheckError();
cudaMemcpy(d_pixel_size,h_pixel_size,sizeof(float),cudaMemcpyHostToDevice);
CudaCheckError();
device_updated = true;
......@@ -867,6 +875,7 @@ void Density::from_device(){
if (device_initiated){
if (padding){
cudaMemcpy(h_data,d_data,h_padding_dim[0]*h_padding_dim[1]*h_padding_dim[2]*sizeof(float),cudaMemcpyDeviceToHost);
CudaCheckError();
} else {
th_data = td_data;
h_data = thrust::raw_pointer_cast(&th_data[0]);
......@@ -1420,9 +1429,10 @@ void create_centric_stamp(float *& d_stamp, int *& d_stamp_offsets, const float
cudaMemcpy(d_stamp_offsets,h_stamp_offsets,sizeof(*h_stamp_offsets)*stamp_vol,cudaMemcpyHostToDevice);
}
Density Density::from_particles_on_grid(Particles &particles, const float &resolution){
void Density::from_particles_on_grid(Density& density,Particles &particles, const float &resolution){
particles.find_bounds();
//
printf("Here4\n");
float * d_stamp;
int * d_stamp_offsets;
uint * d_stamp_pixel_dim;
......@@ -1434,17 +1444,34 @@ Density Density::from_particles_on_grid(Particles &particles, const float &resol
float d;
// //void Sampler::gaussian_topf_extension(float R, float p, const float & h_pixel_size,
// //uint *& d_stamp_pixel_dim, uint *& h_stamp_pixel_dim)
gaussian_topf_extension(resolution, 0.1, h_pixel_size[0], d,d_stamp_pixel_dim, h_stamp_pixel_dim);
// printf("Created stamp with the resolution %f and stamp dimensions %u %u %u\n",resolution,h_stamp_pixel_dim[0],h_stamp_pixel_dim[1],h_stamp_pixel_dim[2]);
gaussian_topf_extension(resolution, 0.1, density.h_pixel_size[0], d,d_stamp_pixel_dim, h_stamp_pixel_dim);
printf("Created stamp with the resolution %f and stamp dimensions %u %u %u\n",resolution,h_stamp_pixel_dim[0],h_stamp_pixel_dim[1],h_stamp_pixel_dim[2]);
//
// //initiate with padding
Density density(particles.h_dimension,h_stamp_pixel_dim,h_pixel_size[0]);
// "/2" since padding_dim is always odd and the "middle pixel" always part of the actual pixel dim
for (int i=0; i<3; i++) density.h_padding_dim[i] = (uint) (density.h_pixel_dim[i] + 2*(h_stamp_pixel_dim[i]/2));
for (int i=0; i<3; i++) density.h_lower_left[i] = 0;
density.padding_index_offset.x = h_stamp_pixel_dim[0]/2;
density.padding_index_offset.y = h_stamp_pixel_dim[1]/2;
density.padding_index_offset.z = h_stamp_pixel_dim[2]/2;
density.padding_linear_offset = density.index_to_linear_space(density.padding_index_offset,density.h_padding_dim);
density.padding = true;
density.th_data.resize(density.h_padding_dim[0]*density.h_padding_dim[1]*density.h_padding_dim[2]);
density.h_data = thrust::raw_pointer_cast(&density.th_data[0]);
density.host_initiated = true;
density.init_device();
CudaCheckError();
// Density density(particles.h_dimension,h_stamp_pixel_dim,h_pixel_size[0]);
CudaCheckError();
density.fill_with(0,true);
particles.center_on_density(density.h_coord_dim);
// //__host__ void create_centric_stamp(float *& d_stamp, uint *& d_stamp_offsets, const float & h_pixel_size,
// // uint * const & h_pixel_dim,uint *& h_stamp_pixel_dim, size_t & stamp_vol,std::function <float (float)> kernel);
create_centric_stamp(d_stamp,d_stamp_offsets,h_pixel_size[0],density.h_padding_dim,h_stamp_pixel_dim,stamp_vol,kernel);
create_centric_stamp(d_stamp,d_stamp_offsets,density.h_pixel_size[0],density.h_padding_dim,h_stamp_pixel_dim,stamp_vol,kernel);
// printf("Centric stamp created\n");
uint * d_offsets;
cudaMalloc((void **) &d_offsets,particles.particle_count()*sizeof(&d_offsets));
......@@ -1457,22 +1484,24 @@ Density Density::from_particles_on_grid(Particles &particles, const float &resol
coords_to_offsets<<<coords_to_offset_grid_dim,coords_to_offset_block_dim>>>
(particles.d_data,d_offsets, particles.d_dimension,density.d_pixel_dim,density.d_padding_dim,
density.padding_linear_offset,particles.particle_count());
//
CudaCheckError();
cudaDeviceSynchronize();
size_t potatoe_stamp_block_dim = 128;
size_t potatoe_stamp_grid_dim = 24;
size_t potatoe_stamp_shared_mem = stamp_vol*sizeof(int) + stamp_vol*sizeof(float);
printf("Potatoe stamp parameters block dim: %u\n",potatoe_stamp_block_dim);
printf("Potatoe stamp parameters grid dim: %u\n",potatoe_stamp_grid_dim);
printf("Potatoe stamp parameters shared mem: %u\n",potatoe_stamp_shared_mem);
//void potatoe_stamp(float4 * d_particles, float * d_density, uint * d_offsets, uint * d_stamp_offsets,
// float * d_stamp, size_t stamp_vol, size_t particle_num){
CudaCheckError();
// potatoe_stamp<<<potatoe_stamp_grid_dim,potatoe_stamp_block_dim,potatoe_stamp_shared_mem>>>
potatoe_stamp<<<potatoe_stamp_grid_dim,potatoe_stamp_block_dim,potatoe_stamp_shared_mem>>>
(particles.d_data, density.d_data ,d_offsets, d_stamp_offsets, d_stamp, stamp_vol, particles.particle_count());
//
CudaCheckError();
cudaDeviceSynchronize();
// auto finish = std::chrono::high_resolution_clock::now();
// std::chrono::duration<double> elapsed = finish - start;
// std::cout << "Elapsed time: " << elapsed.count() << " s\n";
return density;
printf("Done\n");
}
//consumes h_pixel_vol() space
......
......@@ -81,7 +81,7 @@ public:
__host__ void change_canvas(uint * h_new_pixel_dim, uint * old_in_new_offset);
__host__ void mid_cube_with_surface(const uint3& cube_pixel_dim, const uint& edge_thickness,const float& rho_inner,const float& rho_edge);
__host__ std::pair<float3,float3> get_bounding_box(void);
__host__ Density from_particles_on_grid(Particles &particles, const float &resolution);
__host__ static void from_particles_on_grid(Density& density,Particles &particles, const float &resolution);
__host__ Density density_limits_above(float thresold,uint3 tolerance = {2,2,2});
__host__ void difference_map_density_threshold(Density & difference_map,float threshold0,float threshold1, thrust::device_vector<uint>& td_on_surface);
__host__ Density only_linear_indeces(thrust::device_vector<uint>& td_on_surface);
......
......@@ -20,7 +20,7 @@ eval: GpuDelaunay.o KerDivision.o Splaying.o Star.o PredWrapper.o RandGen.o Inpu
ccl: CMrcReader.o
# $(NVCC) --gpu-architecture=sm_61 --include-path=./ $(CUB_INCLUDE) --device-c Labeler.cu Density.cu Kernels.cu Particles.cu TransformationGrid.cu CRotationGrid.cu ccl_test.cu -g -G
$(NVCC) --gpu-architecture=sm_61 --include-path=./ $(CUB_INCLUDE) --device-c TransformationGrid.cu ccl_test.cu -g -G
$(NVCC) --gpu-architecture=sm_61 --include-path=./ $(CUB_INCLUDE) --device-c Density.cu TransformationGrid.cu ccl_test.cu -g -G
$(NVCC) --gpu-architecture=sm_61 --device-link Density.o ccl_test.o Kernels.o Labeler.o Particles.o PdbReader.o TransformationGrid.o CRotationGrid.o --output-file link.o
g++ Density.o Particles.o PdbReader.o CMrcReader.o Kernels.o TransformationGrid.o Labeler.o CRotationGrid.o ccl_test.o link.o -o ccl_test -L. $(CUDA_LINKS) $(BOOST_LINKS) -lpng -lnvidia-ml
......
No preview for this file type
This diff is collapsed.
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