Commit 49ca5467 authored by karius's avatar karius

further steps for alignment protocol

parent f52a8b1d
...@@ -8,7 +8,6 @@ ...@@ -8,7 +8,6 @@
#include <Engine.h> #include <Engine.h>
#include <string> #include <string>
#include <iostream> #include <iostream>
#include <Node.h>
Engine::Engine() { Engine::Engine() {
_workerthreads.resize(_gpumanager.numGPUS()); _workerthreads.resize(_gpumanager.numGPUS());
......
...@@ -6,74 +6,13 @@ ...@@ -6,74 +6,13 @@
*/ */
#include <FittingProtocol.h> #include <FittingProtocol.h>
void FittingProtocol::_determine_gpu_resources(void){
//memory management library
nvmlInit_v2();
cudaGetDeviceCount(&_num_devices);
_device_handles.resize(_num_devices);
_device_memory_infos.resize(_num_devices);
for (int i=0;i<_num_devices;i++){
nvmlDeviceGetHandleByIndex_v2(i,&_device_handles[i]);
nvmlDeviceGetMemoryInfo(_device_handles[i], &_device_memory_infos[i]);
}
CudaCheckError();
}
FittingProtocol::FittingProtocol() { FittingProtocol::FittingProtocol() {}
_determine_gpu_resources();
}
void FittingProtocol::addQueryFromPdb(const char * file_name){
if (file_exists(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()){
ret += target.loadToGpu(gpu_index);
if (ret == 0){
std::cout << "Target has been loaded to device of index " << gpu_index << "..." << std::endl;
}
}
}
else{
std::cout << "Gpu Index " << gpu_index << " does not match the number of devices ..." << std::endl;
}
}
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("-------------------------------------\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("-------------------------------------\n");
}
}
FittingProtocol::~FittingProtocol() { FittingProtocol::~FittingProtocol() {
// TODO Auto-generated destructor stub // TODO Auto-generated destructor stub
} }
void FittingProtocol::setFileTarget(const char * file_name){
if (file_exists(file_name)){
target = Target::fromMrc(file_name);
}
}
...@@ -8,8 +8,6 @@ ...@@ -8,8 +8,6 @@
#ifndef FITTINGPROTOCOL_H_ #ifndef FITTINGPROTOCOL_H_
#define FITTINGPROTOCOL_H_ #define FITTINGPROTOCOL_H_
#include <Target.h>
#include <Query.h>
//#include <Score.h> //#include <Score.h>
#include <vector> #include <vector>
#include <MemoryResources.h> #include <MemoryResources.h>
...@@ -17,19 +15,14 @@ ...@@ -17,19 +15,14 @@
#include <GPUManager.h> #include <GPUManager.h>
#include <vector> #include <vector>
#include <iostream> #include <iostream>
#include <Scoring.h>
class FittingProtocol { class FittingProtocol {
public: public:
FittingProtocol(); FittingProtocol();
void setFileTarget(const char * file_name); void generateScorings(std::vector<Scoring> &scorings);
void addQueryFromPdb(const char * file_name); static void run();
virtual ~FittingProtocol(); virtual ~FittingProtocol();
void printGPUInfo(void);
Target target;
void loadTarget(int gpu_index);
void loadQueries(int gpu_index);
std::vector<Query> queries;
// std::vector<Score> scores;
private: private:
GPUManager _manager; GPUManager _manager;
}; };
......
No preview for this file type
No preview for this file type
...@@ -2,8 +2,11 @@ CUDA_INCLUDE = -I./cuda_util_include #most of it comes with nvcc, this is a conv ...@@ -2,8 +2,11 @@ CUDA_INCLUDE = -I./cuda_util_include #most of it comes with nvcc, this is a conv
BOOST_INCLUDE = -I/usr/local/include/ BOOST_INCLUDE = -I/usr/local/include/
DELAUNAY_INCLUDE = -I./gDel3D/GPU/ DELAUNAY_INCLUDE = -I./gDel3D/GPU/
CUB_INCLUDE = -I/data/tools/cub-1.8.0/ CUB_INCLUDE = -I/data/tools/cub-1.8.0/
EIGEN_INCLUDE = -I/data/tools/lorm/3rdparty/include/eigen3/ #EIGEN_INCLUDE = -I/data/tools/lorm/3rdparty/include/eigen3/
CFLAGS=-I. -dc $(CUDA_INCLUDE) -g -G EIGEN_INCLUDE = -I/data/tools/eigen-eigen-323c052e1731/
CFLAGS=-I. -dc $(CUDA_INCLUDE) -g -G
IMP_INCLUDE = -I/data/git/imp/bin/include/ -I/data/git/imp/build/src/dependency/RMF/include/
IMP_LINKS = -L/data/git/imp/build/lib/ -lRMF -limp_core -limp_rmf -limp_kernel -limp_atom -limp_algebra
BOOST_LINKS = -L/usr/local/lib -lboost_unit_test_framework BOOST_LINKS = -L/usr/local/lib -lboost_unit_test_framework
CUDA_LINKS = -L/usr/local/cuda-10.0/lib64/ -lcusolver -lcudart -lcublas -lcuda -lcudadevrt CUDA_LINKS = -L/usr/local/cuda-10.0/lib64/ -lcusolver -lcudart -lcublas -lcuda -lcudadevrt
...@@ -17,6 +20,11 @@ ifndef DEVICE_ARCH ...@@ -17,6 +20,11 @@ ifndef DEVICE_ARCH
DEVICE_ARCH = sm_61 DEVICE_ARCH = sm_61
endif endif
align:
nvcc --gpu-architecture=$(DEVICE_ARCH) --include-path=./ $(IMP_INCLUDE) $(EIGEN_INCLUDE) --device-c AlignmentProtocol.cu test.cu -g -G
nvcc --gpu-architecture=$(DEVICE_ARCH) --device-link AlignmentProtocol.o test.o --output-file link.o
g++ AlignmentProtocol.o test.o link.o -o align -lnvidia-ml -L/usr/local/cuda-10.0/lib64/ -lcudart -lcudadevrt -pthread $(IMP_LINK)
eval: GpuDelaunay.o KerDivision.o Splaying.o Star.o PredWrapper.o RandGen.o InputCreator.o predicates.o ThrustWrapper.o KerPredicates.o CRotationGrid.o DelaunayChecker.o eval: GpuDelaunay.o KerDivision.o Splaying.o Star.o PredWrapper.o RandGen.o InputCreator.o predicates.o ThrustWrapper.o KerPredicates.o CRotationGrid.o DelaunayChecker.o
nvcc --gpu-architecture=sm_61 --include-path=./,./gDel3D/GPU/ CRotationGrid.o InputCreator.o RandGen.o DelaunayChecker.o KerPredicates.o Splaying.o Star.o predicates.o ThrustWrapper.o PredWrapper.o KerDivision.o GpuDelaunay.o TransformationGrid.cu -o eval nvcc --gpu-architecture=sm_61 --include-path=./,./gDel3D/GPU/ CRotationGrid.o InputCreator.o RandGen.o DelaunayChecker.o KerPredicates.o Splaying.o Star.o predicates.o ThrustWrapper.o PredWrapper.o KerDivision.o GpuDelaunay.o TransformationGrid.cu -o eval
...@@ -33,9 +41,9 @@ manager: ...@@ -33,9 +41,9 @@ manager:
g++ GPUManager.o GPUObject.o Density.o CMrcReader.o test.o link.o -o manager -lnvidia-ml -L/usr/local/cuda-10.0/lib64/ -lcudart -lcudadevrt g++ GPUManager.o GPUObject.o Density.o CMrcReader.o test.o link.o -o manager -lnvidia-ml -L/usr/local/cuda-10.0/lib64/ -lcudart -lcudadevrt
engine: engine:
nvcc --gpu-architecture=$(DEVICE_ARCH) --include-path=./ --device-c Node.cu Parameter.cu ParameterSet.cu ParameterPaths.cu NodeLoadPdb.cu Engine.cu Task.cu WorkerThread.cu GPUManager.cu GPUObject.cu Density.cu CMrcReader.cu test.cu -g -G nvcc --gpu-architecture=$(DEVICE_ARCH) --include-path=./ --device-c ScoreCam.cu ParameterFloat.cu PdbReader.cu Particles.cu ParameterSet.cu ParameterPaths.cu NodeLoadMrc.cu NodeDensityThreshold.cu NodeMolecularDensity.cu NodeLoadPdb.cu Engine.cu Task.cu WorkerThread.cu GPUManager.cu GPUObject.cu Density.cu CMrcReader.cu test.cu -g -G
nvcc --gpu-architecture=$(DEVICE_ARCH) --device-link Node.o Parameter.o ParameterSet.o ParameterPaths.o NodeLoadPdb.o Engine.o Task.o WorkerThread.o GPUManager.o GPUObject.o Density.o CMrcReader.o test.o --output-file link.o nvcc --gpu-architecture=$(DEVICE_ARCH) --device-link ScoreCam.o NodeDensityThreshold.o ParameterFloat.o PdbReader.o Particles.o ParameterSet.o ParameterPaths.o NodeLoadMrc.o NodeLoadPdb.o NodeMolecularDensity.o Engine.o Task.o WorkerThread.o GPUManager.o GPUObject.o Density.o CMrcReader.o test.o --output-file link.o
g++ Node.o Parameter.o ParameterSet.o ParameterPaths.o NodeLoadPdb.o Engine.o Task.o WorkerThread.o GPUManager.o GPUObject.o Density.o CMrcReader.o test.o link.o -o engine -lnvidia-ml -L/usr/local/cuda-10.0/lib64/ -lcudart -lcudadevrt -pthread g++ ScoreCam.o NodeDensityThreshold.o ParameterFloat.o PdbReader.o Particles.o ParameterSet.o ParameterPaths.o NodeLoadMrc.o NodeLoadPdb.o NodeMolecularDensity.o Engine.o Task.o WorkerThread.o GPUManager.o GPUObject.o Density.o CMrcReader.o test.o link.o -o engine -lnvidia-ml -L/usr/local/cuda-10.0/lib64/ -lcudart -lcudadevrt -pthread
protocol: protocol:
# sm_72 volta support, V100s # sm_72 volta support, V100s
......
/*
* Node.cu
*
* Created on: Dec 17, 2020
* Author: kkarius
*/
/* /*
* Node.h * Node.h
* *
* Created on: Dec 17, 2020 * Created on: Dec 21, 2020
* Author: kkarius * Author: kkarius
*/ */
#ifndef NODE_H_ #ifndef NODE_H_
#define NODE_H_ #define NODE_H_
#include <ParameterSet.h>
#include <GPUObject.h> #include <GPUObject.h>
#include <GPUManager.h>
#include <vector>
#include <Engine.h> #include <Engine.h>
#include <ParameterSet.h>
#include <Slot.h> #include <Slot.h>
template <class IN, class OUT> template <class INPUT, class OUTPUT>
class Node : public GPUObject{ class Node : public GPUObject {
public: public:
Node(){}; Node(Engine * _engine){
~Node() {}; this->_manager = &_engine->_gpumanager;
}
ParameterSet * own_paras = nullptr; ParameterSet * own_paras = nullptr;
ParameterSet memory; static void requestParameters(Node * node,ParameterSet * para_set){
void processRequest(const ParameterSet &para_set); node->own_paras->print();
void setGPUManager(GPUManager * manager){ printf("Precursor accessible: %s\n",node->precursor != nullptr ? "True": "False");
this->_manager = manager; if (node->precursor != nullptr)
if (input != nullptr){ requestParameters(node->precursor,para_set);
input->_manager = manager; if (node->own_paras != nullptr)
} para_set->addParameterSet(node->own_paras);
if (output != nullptr){
output->_manager = manager;
}
} }
Slot<IN> * input = nullptr; ParameterSet memory;
Slot<OUT> * output = nullptr; //this shitty implementation is just for the compiler
OUTPUT * request(ParameterSet &para_set){return nullptr;};
Slot<INPUT> input;
Slot<OUTPUT> output;
// OUTPUT * output = new OUTPUT;
// void setPrecursor(void * _precursor){precursor = _precursor;};
// void * precursor = nullptr;
~Node(){};
}; };
#endif /* NODE_H_ */ #endif /* NODE_H_ */
...@@ -10,17 +10,21 @@ ...@@ -10,17 +10,21 @@
#include <ParameterPaths.h> #include <ParameterPaths.h>
#include <MemoryResources.h> #include <MemoryResources.h>
#include <Particles.h> #include <Particles.h>
#include <Node.h> #include <PdbReader.h>
#include <Slot.h>
NodeLoadPdb::NodeLoadPdb(Engine * _engine){ NodeLoadPdb::NodeLoadPdb(Engine * _engine): Node(_engine){
ParameterSet * set_para = new ParameterSet(); ParameterSet * set_para = new ParameterSet;
ParameterPaths * para_paths = new ParameterPaths("pdb_paths"); ParameterPaths * para_paths = new ParameterPaths("pdb_paths");
set_para->addParameter(para_paths); set_para->addParameter(para_paths);
own_paras = set_para; own_paras = set_para;
output = new Slot<Particles>(true); }
setGPUManager(&_engine->_gpumanager); Particles * NodeLoadPdb::request(ParameterSet &para_set){
if (para_set.containsParameter("pdb_paths")){
ParameterPaths * parameter = dynamic_cast<ParameterPaths*>(para_set.getParameter("pdb_paths"));
// Particles::from_pdb(output,parameter->getPath());
}
return nullptr;
} }
NodeLoadPdb::~NodeLoadPdb() { NodeLoadPdb::~NodeLoadPdb() {
......
...@@ -11,11 +11,13 @@ ...@@ -11,11 +11,13 @@
#include <GPUManager.h> #include <GPUManager.h>
#include <Engine.h> #include <Engine.h>
#include <Particles.h> #include <Particles.h>
#include <ParameterSet.h>
#include <Node.h> #include <Node.h>
class NodeLoadPdb: public Node<void,Particles>{ class NodeLoadPdb: public Node<void,Particles>{
public: public:
NodeLoadPdb(Engine * _engine); NodeLoadPdb(Engine * _engine);
Particles * request(ParameterSet &para_set);
virtual ~NodeLoadPdb(); virtual ~NodeLoadPdb();
}; };
......
...@@ -4,10 +4,3 @@ ...@@ -4,10 +4,3 @@
* Created on: Nov 27, 2020 * Created on: Nov 27, 2020
* Author: kkarius * Author: kkarius
*/ */
#include <Parameter.h>
Parameter::~Parameter() {
// TODO Auto-generated destructor stub
}
...@@ -15,8 +15,9 @@ public: ...@@ -15,8 +15,9 @@ public:
Parameter(const char * name):_name(name){}; Parameter(const char * name):_name(name){};
// Parameter(const char * name,float from, float to, float init):_name(name),_from(from),_to(to),_init(init){}; // Parameter(const char * name,float from, float to, float init):_name(name),_from(from),_to(to),_init(init){};
std::string name(void){return std::string(_name);} std::string name(void){return std::string(_name);}
virtual int getVolume(void) = 0;
// float continuousValue(int i = 0){return _init;}; // float continuousValue(int i = 0){return _init;};
virtual ~Parameter(); virtual ~Parameter(){};
private: private:
const char * _name = nullptr; const char * _name = nullptr;
// float _from = 0; // float _from = 0;
......
...@@ -6,8 +6,19 @@ ...@@ -6,8 +6,19 @@
*/ */
#include <ParameterPaths.h> #include <ParameterPaths.h>
#include <util.h>
#include <algorithm>
ParameterPaths::~ParameterPaths() { ParameterPaths::~ParameterPaths() {
// TODO Auto-generated destructor stub // TODO Auto-generated destructor stub
} }
void ParameterPaths::addPath(const char * path){
if (file_exists(path) && _paths.end() == std::find(_paths.begin(),_paths.end(),path)){
_paths.push_back(path);
}
}
int ParameterPaths::getVolume(){
return _paths.size();
}
...@@ -13,7 +13,10 @@ ...@@ -13,7 +13,10 @@
class ParameterPaths : public Parameter { class ParameterPaths : public Parameter {
public: public:
ParameterPaths(const char * _name):Parameter(_name){}; ParameterPaths(const char * name):Parameter(name){};
const char * getPath(int index = 0){return _paths[index];}
void addPath(const char * path);
int getVolume(void);
virtual ~ParameterPaths(); virtual ~ParameterPaths();
private: private:
std::vector<const char *> _paths; std::vector<const char *> _paths;
......
...@@ -6,13 +6,47 @@ ...@@ -6,13 +6,47 @@
*/ */
#include <ParameterSet.h> #include <ParameterSet.h>
#include <algorithm>
ParameterSet::ParameterSet() { int ParameterSet::getVolume(void){
// TODO Auto-generated constructor stub int ret = 1;
for (auto const &par:_parameters){
ret *= par.second->getVolume();
}
for (ParameterSet * parset:_parameter_sets){
ret *= parset->getVolume();
}
return ret;
} }
ParameterSet::~ParameterSet() { int ParameterSet::getMultiIndexSize(void){
// TODO Auto-generated destructor stub return _parameters.size() + _parameter_sets.size();
} }
std::vector<int> ParameterSet::getMultiIndex(int i){
std::vector<int> ret(_parameters.size() + _parameter_sets.size());
int p = _parameters.size();
int ps = _parameter_sets.size();
int d = p + ps;
std::vector<int> volumes(d);
for (int j = 0;j<ps;j++){
volumes[j] = _parameter_sets[j]->getVolume();
}
for (int j = 0;j<p;j++){
volumes[ps+j] = _parameter_sets[j]->getVolume();
}
int l = 0;
int v;
int e;
for (int j = 0;j<d-1;j++){
v = 1;
for (int k=0;k<d-j;k++){
v *= volumes[k];
}
e = (i-l)/v;
ret[j] = e;
l += e*v;
}
ret[d-1] = i-l;
return ret;
}
...@@ -11,15 +11,33 @@ ...@@ -11,15 +11,33 @@
#include <Parameter.h> #include <Parameter.h>
#include <map> #include <map>
#include <string> #include <string>
#include <vector>
class ParameterSet { class ParameterSet {
public: public:
ParameterSet(); ParameterSet(){};
void addParameter(Parameter * parameter){_parameters[parameter->name()] = parameter;} void addParameter(Parameter * parameter){_parameters[parameter->name()] = parameter;}
void addParameterSet(ParameterSet * parameter_set){_parameter_sets.push_back(parameter_set);};
Parameter * getParameter(const char * name){return _parameters[std::string(name)];} Parameter * getParameter(const char * name){return _parameters[std::string(name)];}
virtual ~ParameterSet(); bool containsParameter(const char * name){return bool(_parameters.end() != _parameters.find(name));}
int getVolume(void);
int getMultiIndexSize(void);
std::vector<int> getMultiIndex(int i);
void print(void){
printf("========\n");
printf("Parameters:\n");
for (auto const& para: _parameters){
printf("%s\n",para.first.c_str());
}
printf("Parameter Sets:\n");
for (auto const& para_set:_parameter_sets){
para_set->print();
}
}
virtual ~ParameterSet(){};
private: private:
std::map<std::string,Parameter *> _parameters; std::map<std::string,Parameter *> _parameters;
std::vector<ParameterSet *> _parameter_sets;
}; };
#endif /* PARAMETERSET_H_ */ #endif /* PARAMETERSET_H_ */
...@@ -8,21 +8,29 @@ ...@@ -8,21 +8,29 @@
#include <Particles.h> #include <Particles.h>
Particles::Particles(){ Particles::Particles(){
int error = 0; // bool gpu_ok = true;
h_bounding_box = (float *)malloc(6*sizeof(*h_bounding_box)); // h_bounding_box = (float *)malloc(6*sizeof(*h_bounding_box));
error = int(cudaSuccess != cudaMalloc((void **) &d_bounding_box, 6*sizeof(*h_bounding_box))); // gpu_ok &= blockMemory(this,gpu_index,d_bounding_box, 6*sizeof(*d_bounding_box),h_bounding_box,true);
h_dimension = (float *) malloc(3*sizeof(*h_dimension)); // h_dimension = (float *) malloc(3*sizeof(*h_dimension));
error = int(cudaSuccess != cudaMalloc((void **) &d_dimension,3*sizeof(*d_dimension))); // gpu_ok &= blockMemory(this,gpu_index,d_dimension, 6*sizeof(*d_dimension),h_dimension,true);
error = int(cudaSuccess != cudaMalloc((void **) &d_particle_count,sizeof(*d_particle_count))); // gpu_ok &= blockMemory(this,gpu_index,d_particle_count, sizeof(*d_particle_count),h_particle_count,true);
assert(error == 0); // assert(gpu_ok == true);
} }
Particles::~Particles() {} bool Particles::createInstance(int gpu_index){
Particles_Register instance;
bool gpu_ok = true;
instance.h_bounding_box = (float *)malloc(6*sizeof(*instance.h_bounding_box));
gpu_ok &= blockMemory(this,gpu_index,instance.d_bounding_box, 6*sizeof(*instance.d_bounding_box),instance.h_bounding_box,true);
instance.h_dimension = (float *) malloc(3*sizeof(*instance.h_dimension));
gpu_ok &= blockMemory(this,gpu_index,instance.d_dimension, 6*sizeof(*instance.d_dimension),instance.h_dimension,true);
gpu_ok &= blockMemory(this,gpu_index,instance.d_particle_count, sizeof(*instance.d_particle_count),instance.h_particle_count,true);
assert(gpu_ok == true);
_instance_registry[gpu_index] = instance;
return gpu_ok;
}
void Particles::set_particle_count(const size_t & count){ Particles::~Particles() {}
m_particle_count = count;
cudaMemcpy(d_particle_count,&m_particle_count,sizeof(size_t),cudaMemcpyHostToDevice);
};
struct translate : public thrust::unary_function<float4,float4> struct translate : public thrust::unary_function<float4,float4>
{ {
...@@ -33,80 +41,68 @@ struct translate : public thrust::unary_function<float4,float4> ...@@ -33,80 +41,68 @@ struct translate : public thrust::unary_function<float4,float4>
} }
}; };
//__host__
//void Particles::center_on_density(float * const& h_coord_dim){
// find_bounds();
// 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(thrust::device_pointer_cast(&d_data[0]),thrust::device_pointer_cast(&d_data[h_particle_count[0]-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__ //__host__
void Particles::center_on_density(float * const& h_coord_dim){ //void Particles::origin_to_center(){
find_bounds(); // find_bounds();
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}; // 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; // translate unary_function_shift;
unary_function_shift.shift = shift; // unary_function_shift.shift = shift;
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); // thrust::transform(thrust::device_pointer_cast(&d_data[0]),thrust::device_pointer_cast(&d_data[h_particle_count[0]-1]),thrust::device_pointer_cast(&d_data[0]),unary_function_shift);
h_bounding_box[0] += shift.x; // h_bounding_box[0] += shift.x;
h_bounding_box[1] += shift.y; // h_bounding_box[1] += shift.y;
h_bounding_box[2] += shift.z; // h_bounding_box[2] += shift.z;
h_bounding_box[3] += shift.x;