Skip to content

GitLab

Commit 849f588e authored by karius's avatar karius
Browse files
Options

added first queries

parent ad047d6b
Hide whitespace changes
Inline Side-by-side
Showing with 203 additions and 133 deletions
Density.cu
View file @ 849f588e
......@@ -576,20 +576,20 @@ std::pair<float3,float3> Density::get_bounding_box(void){
return ret;
}
__host__
Particles Density::particles_from_density(const thrust::device_vector<uint>& td_linear_indices){
Particles ret;
ret.td_data.resize(td_linear_indices.size());
ret.set_particle_count(td_linear_indices.size());
linear_offset_to_coord unary_transform_function;
unary_transform_function.d_pixel_dim = d_pixel_dim;
unary_transform_function.pixel_size = h_pixel_size[0];
thrust::transform(td_linear_indices.begin(),td_linear_indices.end(),ret.td_data.begin(),unary_transform_function);
ret.th_data = ret.td_data;
ret.d_data = thrust::raw_pointer_cast(&ret.td_data[0]);
ret.h_data = thrust::raw_pointer_cast(&ret.th_data[0]);
return ret;
}
//__host__
//Particles Density::particles_from_density(const thrust::device_vector<uint>& td_linear_indices){
// Particles ret;
// ret.td_data.resize(td_linear_indices.size());
// ret.set_particle_count(td_linear_indices.size());
// linear_offset_to_coord unary_transform_function;
// unary_transform_function.d_pixel_dim = d_pixel_dim;
// unary_transform_function.pixel_size = h_pixel_size[0];
// thrust::transform(td_linear_indices.begin(),td_linear_indices.end(),ret.td_data.begin(),unary_transform_function);
// ret.th_data = ret.td_data;
// ret.d_data = thrust::raw_pointer_cast(&ret.td_data[0]);
// ret.h_data = thrust::raw_pointer_cast(&ret.th_data[0]);
// return ret;
//}
__host__
void Density::change_canvas(uint * h_new_pixel_dim, uint * old_in_new_offset){
......
Density.h
View file @ 849f588e
......@@ -85,7 +85,7 @@ public:
__host__ Density only_linear_indeces(thrust::device_vector<uint>& td_on_surface);
__host__ void to_mrc_write_only(thrust::device_vector<uint> & td_linear_indices,const uint & label, const char * mrc_path, uint log_2_dim);
__host__ void cut_and_binarize(Density & sub_density,float threshold,float above_value,float below_value,uint3 tolerance);
__host__ Particles particles_from_density(const thrust::device_vector<uint>& td_linear_indices);
// __host__ Particles particles_from_density(const thrust::device_vector<uint>& td_linear_indices);
__host__ uint3 linear_to_index_space(const uint &linear_index, uint * const &pixel_dim);
__host__ uint index_to_linear_space(const uint3 &pixel_index, uint * const &pixel_dim);
//__host__ uint linear_to_linear_space(const uint &linear_pixel_index, uint * const &pixel_dimA, uint * const &pixel_dimB,uint3 relativ_offset);
......
FittingProtocol.cu
View file @ 849f588e
......@@ -23,10 +23,19 @@ FittingProtocol::FittingProtocol() {
_determine_gpu_resources();
}
void FittingProtocol::addQueryFromPdb(const char * file_name){
Query query = Query::fromPdb(file_name);
queries.push_back(query);
}
void FittingProtocol::loadTarget(int gpu_index){
int ret = 0;
if(gpu_index < _num_devices){
if(_device_memory_infos[gpu_index].free > target.getGpuMem()){
target.loadToGpu(gpu_index);
ret += target.loadToGpu(gpu_index);
if (ret == 0){
std::cout << "Target has been loaded to device of index " << gpu_index << "..." << std::endl;
}
}
}
else{
......@@ -34,13 +43,25 @@ void FittingProtocol::loadTarget(int gpu_index){
}
}
void FittingProtocol::loadQueries(int gpu_index){
int error = 0;
for (Query query: queries){
if(_device_memory_infos[gpu_index].free > query.getGpuMem()){
error += query.loadToGpu(gpu_index);
if (error == 0){
std::cout << "Queries have been loaded to device of index " << gpu_index << "..." << std::endl;
}
}
}
}
void FittingProtocol::printGPUInfo(void){
printf("Found %i devices ...\n",_num_devices);
printf("-------------------------------------");
printf("-------------------------------------\n");
for (int i=0;i<_num_devices;i++){
printf("Device %i:\n",i);
printf("Total Memory: %u bytes\n",_device_memory_infos[i].total);
printf("-------------------------------------");
printf("-------------------------------------\n");
}
}
......
FittingProtocol.h
View file @ 849f588e
......@@ -9,7 +9,7 @@
#define FITTINGPROTOCOL_H_
#include <Target.h>
//#include <Query.h>
#include <Query.h>
//#include <Score.h>
#include <vector>
#include <MemoryResources.h>
......@@ -22,11 +22,13 @@ class FittingProtocol {
public:
FittingProtocol();
void setFileTarget(const char * file_name);
void addQueryFromPdb(const char * file_name);
virtual ~FittingProtocol();
void printGPUInfo(void);
Target target;
void loadTarget(int gpu_index);
// std::vector<Query> queries;
void loadQueries(int gpu_index);
std::vector<Query> queries;
// std::vector<Score> scores;
private:
size_t _byte_vol = 0;
......
Makefile.test
View file @ 849f588e
......@@ -33,9 +33,9 @@ protocol:
# nvcc --gpu-architecture=sm_61 --device-link Particles.o PdbReader.o FittingProtocol.o Density.o Target.o CMrcReader.o --output-file link.o -lnvidia-ml $(CUDA_LINKS)
# g++ FittingProtocol.o Particles.o PdbReader.o Target.o Density.o CMrcReader.o ccl_test.o link.o -o protocol -lnvidia-ml -L. $(CUDA_LINKS)
$(info Compiling for device architecture $(DEVICE_ARCH))
nvcc --gpu-architecture=$(DEVICE_ARCH) --include-path=./ --device-c PdbReader.cu Particles.cu CMrcReader.cu Density.cu Target.cu FittingProtocol.cu ccl_test.cu -g -G
nvcc --gpu-architecture=$(DEVICE_ARCH) --device-link PdbReader.o Particles.o CMrcReader.o Density.o Target.o FittingProtocol.o ccl_test.o --output-file link.o
g++ PdbReader.o Particles.o CMrcReader.o Density.o Target.o FittingProtocol.o ccl_test.o link.o -o protocol -lnvidia-ml -L/usr/local/cuda-10.0/lib64/ -lcudart -lcudadevrt
nvcc --gpu-architecture=$(DEVICE_ARCH) --include-path=./ --device-c PdbReader.cu Particles.cu CMrcReader.cu Density.cu Target.cu Query.cu FittingProtocol.cu ccl_test.cu -g -G
nvcc --gpu-architecture=$(DEVICE_ARCH) --device-link PdbReader.o Particles.o CMrcReader.o Density.o Target.o Query.o FittingProtocol.o ccl_test.o --output-file link.o
g++ PdbReader.o Particles.o CMrcReader.o Density.o Target.o FittingProtocol.o Query.o ccl_test.o link.o -o protocol -lnvidia-ml -L/usr/local/cuda-10.0/lib64/ -lcudart -lcudadevrt
ccl2: CMrcReader.o
......
Particles.cu
View file @ 849f588e
......@@ -8,13 +8,13 @@
#include <Particles.h>
Particles::Particles(){
th_bounding_box.resize(6);
td_bounding_box.resize(6);
d_bounding_box = thrust::raw_pointer_cast(&td_bounding_box[0]);
h_dimension = (float *) malloc(3*sizeof(float));
cudaMalloc((void **) &d_dimension,3*sizeof(float));
cudaMalloc((void **) &d_particle_count,sizeof(size_t));
CudaCheckError();
int error = 0;
h_bounding_box = (float *)malloc(6*sizeof(*h_bounding_box));
error = int(cudaSuccess != cudaMalloc((void **) &d_bounding_box, 6*sizeof(*h_bounding_box)));
h_dimension = (float *) malloc(3*sizeof(*h_dimension));
error = int(cudaSuccess != cudaMalloc((void **) &d_dimension,3*sizeof(*d_dimension)));
error = int(cudaSuccess != cudaMalloc((void **) &d_particle_count,sizeof(*d_particle_count)));
assert(error == 0);
}
Particles::~Particles() {}
......@@ -33,92 +33,70 @@ struct translate : public thrust::unary_function<float4,float4>
}
};
__host__
std::pair<float3,float3> Particles::get_bounding_box(void){
std::pair<float3,float3> ret;
ret.first = make_float3(th_bounding_box[0],th_bounding_box[1],th_bounding_box[2]);
ret.second = make_float3(th_bounding_box[3],th_bounding_box[4],th_bounding_box[5]);
return ret;
}
__host__
void Particles::center_on_density(float * const& h_coord_dim){
find_bounds();
float4 shift = {h_coord_dim[0]/2 - (th_bounding_box[0] + h_dimension[0]/2),h_coord_dim[1]/2 - (th_bounding_box[1] + h_dimension[1]/2),h_coord_dim[2]/2 - (th_bounding_box[2] + h_dimension[2]/2),0};
float4 shift = {h_coord_dim[0]/2 - (h_bounding_box[0] + h_dimension[0]/2),h_coord_dim[1]/2 - (h_bounding_box[1] + h_dimension[1]/2),h_coord_dim[2]/2 - (h_bounding_box[2] + h_dimension[2]/2),0};
translate unary_function_shift;
unary_function_shift.shift = shift;
thrust::transform(td_data.begin(),td_data.end(),td_data.begin(),unary_function_shift);
th_bounding_box[0] += shift.x;
th_bounding_box[1] += shift.y;
th_bounding_box[2] += shift.z;
th_bounding_box[3] += shift.x;
th_bounding_box[4] += shift.y;
th_bounding_box[5] += shift.z;
td_bounding_box = th_bounding_box;
thrust::transform(thrust::device_pointer_cast(&d_data[0]),thrust::device_pointer_cast(&d_data[m_particle_count-1]),thrust::device_pointer_cast(&d_data[0]),unary_function_shift);
h_bounding_box[0] += shift.x;
h_bounding_box[1] += shift.y;
h_bounding_box[2] += shift.z;
h_bounding_box[3] += shift.x;
h_bounding_box[4] += shift.y;
h_bounding_box[5] += shift.z;
assert(cudaSuccess == cudaMemcpy(d_bounding_box,h_bounding_box,6*sizeof(*d_bounding_box),cudaMemcpyHostToDevice));
}
__host__
void Particles::origin_to_center(){
find_bounds();
float4 shift = {-(th_bounding_box[3] + th_bounding_box[0])/2,-(th_bounding_box[4] + th_bounding_box[1])/2,-(th_bounding_box[5] + th_bounding_box[2])/2,0};
float4 shift = {-(h_bounding_box[3] + h_bounding_box[0])/2,-(h_bounding_box[4] + h_bounding_box[1])/2,-(h_bounding_box[5] + h_bounding_box[2])/2,0};
translate unary_function_shift;
unary_function_shift.shift = shift;
thrust::transform(td_data.begin(),td_data.end(),td_data.begin(),unary_function_shift);
d_data = thrust::raw_pointer_cast(&td_data[0]);
th_bounding_box[0] += shift.x;
th_bounding_box[1] += shift.y;
th_bounding_box[2] += shift.z;
th_bounding_box[3] += shift.x;
th_bounding_box[4] += shift.y;
th_bounding_box[5] += shift.z;
td_bounding_box = th_bounding_box;
thrust::transform(thrust::device_pointer_cast(&d_data[0]),thrust::device_pointer_cast(&d_data[m_particle_count-1]),thrust::device_pointer_cast(&d_data[0]),unary_function_shift);
h_bounding_box[0] += shift.x;
h_bounding_box[1] += shift.y;
h_bounding_box[2] += shift.z;
h_bounding_box[3] += shift.x;
h_bounding_box[4] += shift.y;
h_bounding_box[5] += shift.z;
assert(cudaSuccess == cudaMemcpy(d_bounding_box,h_bounding_box,6*sizeof(*d_bounding_box),cudaMemcpyHostToDevice));
}
Particles Particles::from_pdb(const char * file_path, const bool & to_device = true){
Particles Particles::from_pdb(const char * file_path){
Particles particles;
PdbReader reader(file_path);
particles.pdb_readers.push_back(&reader);
reader.atom_coords_to_device(particles.td_data);
particles.d_data = thrust::raw_pointer_cast(&particles.td_data[0]);
particles.set_particle_count(reader.get_atom_count());
return particles;
}
void Particles::add_particle(const float4 & coords){
set_particle_count(m_particle_count + 1);
td_data.resize(m_particle_count);
td_data[m_particle_count-1] = coords;
d_data = thrust::raw_pointer_cast(&td_data[0]);
CudaCheckError();
}
void Particles::find_bounds(){
float4 max = {FLT_MAX,FLT_MAX,FLT_MAX,0};
float4 min = {0,0,0,0};
float4 min_coords = thrust::reduce(td_data.begin(),td_data.end(),max,min_float4());
float4 max_coords = thrust::reduce(td_data.begin(),td_data.end(),min,max_float4());
th_bounding_box[0] = min_coords.x;
th_bounding_box[1] = min_coords.y;
th_bounding_box[2] = min_coords.z;
th_bounding_box[3] = max_coords.x;
th_bounding_box[4] = max_coords.y;
th_bounding_box[5] = max_coords.z;
td_bounding_box = th_bounding_box;
float4 min_coords = thrust::reduce(thrust::device_pointer_cast(&d_data[0]),thrust::device_pointer_cast(&d_data[m_particle_count-1]),max,min_float4());
float4 max_coords = thrust::reduce(thrust::device_pointer_cast(&d_data[0]),thrust::device_pointer_cast(&d_data[m_particle_count-1]),min,max_float4());
h_bounding_box[0] = min_coords.x;
h_bounding_box[1] = min_coords.y;
h_bounding_box[2] = min_coords.z;
h_bounding_box[3] = max_coords.x;
h_bounding_box[4] = max_coords.y;
h_bounding_box[5] = max_coords.z;
assert(cudaSuccess == cudaMemcpy(d_bounding_box,h_bounding_box,6*sizeof(*d_bounding_box),cudaMemcpyHostToDevice));
//calc dimension
for (int i=0;i<3;i++)
h_dimension[i] = th_bounding_box[i+3] - th_bounding_box[i];
cudaMemcpy(d_dimension,h_dimension,3*sizeof(float),cudaMemcpyHostToDevice);
CudaCheckError();
h_dimension[i] = h_bounding_box[i+3] - h_bounding_box[i];
assert(cudaSuccess == cudaMemcpy(d_dimension,h_dimension,3*sizeof(float),cudaMemcpyHostToDevice));
}
__host__
std::pair<float3,float3> Particles::get_rotationally_safe_bounding_box(void){
std::pair<float3,float3> Particles::rotationally_safe_bounding_box(void){
std::pair<float3,float3> ret;
float d = sqrtf(powf(th_bounding_box[3]-th_bounding_box[0],2) + powf(th_bounding_box[1]-th_bounding_box[4],2) + powf(th_bounding_box[2]-th_bounding_box[5],2))/2;
float d = sqrtf(powf(h_bounding_box[3]-h_bounding_box[0],2) + powf(h_bounding_box[1]-h_bounding_box[4],2) + powf(h_bounding_box[2]-h_bounding_box[5],2))/2;
ret.first = make_float3(-d,-d,-d);
ret.second = make_float3(d,d,d);
return ret;
......@@ -129,6 +107,6 @@ void Particles::d_get_dimension(float *& d_dim){
}
void Particles::print_bounding_box(){
printf("Bounding box lower left front: %f %f %f\n",th_bounding_box[0],th_bounding_box[1],th_bounding_box[2]);
printf("Bounding box upper right back: %f %f %f\n",th_bounding_box[3],th_bounding_box[4],th_bounding_box[5]);
printf("Bounding box lower left front: %f %f %f\n",h_bounding_box[0],h_bounding_box[1],h_bounding_box[2]);
printf("Bounding box upper right back: %f %f %f\n",h_bounding_box[3],h_bounding_box[4],h_bounding_box[5]);
}
Particles.h
View file @ 849f588e
......@@ -8,54 +8,53 @@
#ifndef PARTICLES_H_
#define PARTICLES_H_
#include <thrust/pair.h>
#include <thrust/copy.h>
#include <thrust/device_malloc.h>
#include <thrust/host_vector.h>
#include <thrust/device_vector.h>
#include <thrust/extrema.h>
#include <thrust/execution_policy.h>
#include <thrust/reduce.h>
#include <thrust/functional.h>
#include <thrust/transform.h>
#include <string>
#include <vector>
#include <cmath>
#include <cstdio>
#include <cctype>
#include <cuda_util_include/cutil_math.h>
#include <cuda_runtime.h>
#include <cuda.h>
#include <cuda_device_runtime_api.h>
#include <sstream>
#include <regex>
#include <iostream>
#include <util.h>
#include <assert.h>
#include <cuda_util_include/cutil_math.h>
#include <cuda_runtime.h>
#include <cuda.h>
#include <cuda_device_runtime_api.h>
#include <thrust/pair.h>
#include <thrust/copy.h>
#include <thrust/device_malloc.h>
#include <thrust/device_ptr.h>
#include <thrust/extrema.h>
#include <thrust/execution_policy.h>
#include <thrust/reduce.h>
#include <thrust/functional.h>
#include <thrust/transform.h>
#include <PdbReader.h>
class Particles {
public:
Particles();
virtual ~Particles();
size_t particle_count(){return m_particle_count;};
void set_particle_count(const size_t & count);
static Particles from_pdb(const char * file_path, const bool & to_device);
static Particles from_pdb(const char * file_path);
std::vector<PdbReader *> pdb_readers;
void find_bounds(void);
void print_bounding_box(void);
void add_particle(const float4 &);
thrust::host_vector<float4> th_data;
thrust::device_vector<float4> td_data;
float4 * h_data;
float4 * d_data;
thrust::host_vector<float> th_bounding_box;
thrust::device_vector<float> td_bounding_box;
float * d_bounding_box;
float * d_dimension;
float * h_dimension;
size_t * d_particle_count;
__host__ std::pair<float3,float3> get_bounding_box(void);
__host__ std::pair<float3,float3> get_rotationally_safe_bounding_box(void);
float4 * h_data = nullptr;
float4 * d_data = nullptr;
float * h_bounding_box= nullptr;
float * d_bounding_box= nullptr;
float * d_dimension = nullptr;
float * h_dimension = nullptr;
size_t * d_particle_count = nullptr;
__host__ std::pair<float3,float3> rotationally_safe_bounding_box(void);
__host__ void center_on_density(float * const& h_coord_dim);
__device__ void d_get_dimension(float *&);
__host__ void origin_to_center(void);
......
PdbReader.cu
View file @ 849f588e
......@@ -43,20 +43,24 @@ float4 PdbReader::get_atom(std::string & line){
return atom_coords;
}
void PdbReader::atom_coords_to_device(thrust::device_vector<float4>& td_coords){
int PdbReader::atom_coords_to_device(float4 * d_coords, int gpu_index){
RecordType type;
std::string line;
ifs.clear();
ifs.seekg(0,std::ios::beg);
td_coords.resize(get_atom_count());
int c = 0;
assert(cudaSuccess == cudaSetDevice(gpu_index));
float4 current_atom;
int error = 0;
while (std::getline(ifs,line)){
type = get_type(line);
if (type == ATOM || type == HETATM){
td_coords[c] = get_atom(line);
current_atom = get_atom(line);
error += int(cudaSuccess == cudaMemcpy(d_coords + c,&current_atom,sizeof(*d_coords),cudaMemcpyHostToDevice));
c++;
}
}
return error;
}
//STOLEN from Chimera
......@@ -338,15 +342,3 @@ RecordType PdbReader::get_type(std::string line) {
return UNKNOWN;
}
//int main(int argc, char * argv []){
// if (argc == 2){
// PdbReader r(argv[1]);
// printf("Atom number:%u\n",r.get_atom_count());
// float4 * d_coords;
// r.get_atom_coords(d_coords,true);
// }
//}
PdbReader.h
View file @ 849f588e
......@@ -16,6 +16,7 @@
#include <iostream>
#include <fstream>
#include <list>
#include <assert.h>
#include <cuda_util_include/cutil_math.h>
#include <thrust/device_vector.h>
......@@ -363,7 +364,7 @@ public:
RecordType get_type(std::string line);
static void INIT_AA_321_MAP(void);
static std::string TRANSLATE_AA_321(std::string three_letter_aa_code);
void atom_coords_to_device(thrust::device_vector<float4> & td_coords);
int atom_coords_to_device(float4 * d_coords, int gpu_index);
private:
void scan_file(std::ifstream & ifs);
AtomType get_atom_type(std::string & line);
......
Query.cu
View file @ 849f588e
......@@ -14,3 +14,66 @@ Query::~Query() {
// TODO Auto-generated destructor stub
}
Query Query::fromPdb(const char * fileName){
Query query;
query._source_type = PDB;
query._particles = Particles::from_pdb(fileName);
return query;
}
size_t Query::getCpuMem(){
if (_source_type == PDB){
return _particles.particle_count()*sizeof(*_particles.h_data);
}
return 0;
}
size_t Query::getGpuMem(){
if (_source_type == PDB){
return _particles.particle_count()*sizeof(*_particles.h_data);
}
return 0;
}
int Query::initCpu(){
int ret = 0;
if (_source_type == PDB){
//doesn't really need cpu representation
}
return ret;
}
int Query::initGpu(int gpu_index){
int ret = 0;
if (_source_type == PDB){
//somewhat odd, not clear why the first pdb reader. sort this out from the use case perspective
ret += int(cudaSuccess != cudaSetDevice(gpu_index));
ret += int(cudaSuccess != cudaMalloc((void **) &_particles.d_data,getGpuMem()));
}
return ret;
}
int Query::toGpu(int gpu_index){
int ret = 0;
if (_source_type == PDB){
ret += int(cudaSuccess != cudaSetDevice(gpu_index));
ret += _particles.pdb_readers[0]->atom_coords_to_device(_particles.d_data,gpu_index);
}
return ret;
}
int Query::loadToGpu(int gpu_index){
int ret = 0;
if (!_cpu_initiated){
ret += initCpu();
}
assert(ret != 0);
if (!_gpu_initiated){
ret += initGpu(gpu_index);
}
assert(ret != 0);
toGpu(gpu_index);
assert(ret != 0);
return ret;
}
int Query::fromGpu(){return 0;}
Query.h
View file @ 849f588e
......@@ -19,12 +19,19 @@ public:
size_t getCpuMem();
size_t getGpuMem();
int initCpu();
int initGpu();
int toGpu();
int initGpu(int gpu_index);
int toGpu(int gpu_index);
int fromGpu();
int loadToGpu(int gpu_index);
virtual ~Query();
// Density * as_density();
// Particles * as_particles();
private:
SourceType _source_type = NONE;
bool _cpu_initiated = true; //not really needed currently
bool _gpu_initiated = false;
Particles _particles;
Density _density;
};
#endif /* QUERY_H_ */
Target.cu
View file @ 849f588e
......@@ -26,10 +26,13 @@ int Target::loadToGpu(int gpu_index){
if (!_cpu_initiated){
ret += initCpu();
}
assert(ret != 0);
if (!_gpu_initiated){
ret += initGpu(gpu_index);
}
toGpu(gpu_index);
assert(ret != 0);
ret += toGpu(gpu_index);
assert(ret != 0);
return ret;
}
......
ccl_test.cu
View file @ 849f588e
......@@ -34,6 +34,10 @@ BOOST_AUTO_TEST_CASE(fitting_protocol)
printf(" ============================= ++++++++++++++++++++++++++++++++ ===========================\n");
FittingProtocol protocol;
protocol.printGPUInfo();
protocol.setFileTarget("test/TauA_3.47A_map.mrc");
protocol.loadTarget(0);
protocol.addQueryFromPdb("test/TauA.pdb");
protocol.loadQueries(0);
// std::pair<float3,float3> effective_translational_shift;
// effective_translational_shift.first = {0,0,0};
// effective_translational_shift.second = {2,2,2};
......
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