Skip to content

GitLab

B
BioEM
Commit 47c12f7a authored by David Rohr's avatar David Rohr
Browse files
Options

Merge FFT Algorithm into BioEM version in git repository, cleanup, ensure... 

Merge FFT Algorithm into BioEM version in git repository, cleanup, ensure consistent indentation style
parent 73b7953e
Expand all Hide whitespace changes
Inline Side-by-side
Showing with 1359 additions and 1090 deletions
bioem.cpp
View file @ 47c12f7a
This diff is collapsed.
bioem_algorithm.h
View file @ 47c12f7a
......@@ -15,22 +15,22 @@ template <int GPUAlgo, class RefT>
__device__ static inline void compareRefMap(const int iRefMap, const int iOrient, const int iConv, const bioem_map& Mapconv, bioem_Probability* pProb, const bioem_param_device& param, const RefT& RefMap,
const int cent_x, const int cent_y, const int myShift = 0, const int nShifts2 = 0, const int myRef = 0, const bool threadActive = true)
{
/**************************************************************************************/
/********************** Calculating BioEM Probability ********************************/
/************************* Loop of center displacement here ***************************/
/**************************************************************************************/
/********************** Calculating BioEM Probability ********************************/
/************************* Loop of center displacement here ***************************/
// Taking into account the center displacement
// Taking into account the center displacement
/*** Inizialzing crosscorrelations of calculated projected convolutions ***/
/*** Inizialzing crosscorrelations of calculated projected convolutions ***/
#ifdef SSECODE
myfloat_t sum, sumsquare, crossproMapConv;
__m128 sum_v = _mm_setzero_ps(), sumsquare_v = _mm_setzero_ps(), cross_v = _mm_setzero_ps(), d1, d2;
#else
myfloat_t sum=0.0;
myfloat_t sumsquare=0.0;
myfloat_t crossproMapConv=0.0;
myfloat_t sumsquare=0.0;
myfloat_t crossproMapConv=0.0;
#endif
/****** Loop over Pixels to calculate dot product and cross-correlations of displaced Ref Conv. Map***/
/****** Loop over Pixels to calculate dot product and cross-correlations of displaced Ref Conv. Map***/
myfloat_t logpro;
if (GPUAlgo != 2 || threadActive)
{
......@@ -287,47 +287,47 @@ __device__ static inline void compareRefMap(const int iRefMap, const int iOrient
else
#endif
/***** Summing & Storing total/Orientation Probabilites for each map ************/
{
/******* Summing total Probabilities *************/
/******* Need a constant because of numerical divergence*****/
if(pProb[iRefMap].Constoadd < logpro)
{
pProb[iRefMap].Total = pProb[iRefMap].Total * exp(-logpro + pProb[iRefMap].Constoadd);
pProb[iRefMap].Constoadd = logpro;
}
pProb[iRefMap].Total += exp(logpro - pProb[iRefMap].Constoadd);
/***** Summing & Storing total/Orientation Probabilites for each map ************/
{
/******* Summing total Probabilities *************/
/******* Need a constant because of numerical divergence*****/
if(pProb[iRefMap].Constoadd < logpro)
{
pProb[iRefMap].Total = pProb[iRefMap].Total * exp(-logpro + pProb[iRefMap].Constoadd);
pProb[iRefMap].Constoadd = logpro;
}
pProb[iRefMap].Total += exp(logpro - pProb[iRefMap].Constoadd);
//Summing probabilities for each orientation
if(pProb[iRefMap].ConstAngle[iOrient] < logpro)
{
pProb[iRefMap].forAngles[iOrient] = pProb[iRefMap].forAngles[iOrient] * exp(-logpro + pProb[iRefMap].ConstAngle[iOrient]);
pProb[iRefMap].ConstAngle[iOrient] = logpro;
if(pProb[iRefMap].ConstAngle[iOrient] < logpro)
{
pProb[iRefMap].forAngles[iOrient] = pProb[iRefMap].forAngles[iOrient] * exp(-logpro + pProb[iRefMap].ConstAngle[iOrient]);
pProb[iRefMap].ConstAngle[iOrient] = logpro;
}
pProb[iRefMap].forAngles[iOrient] += exp(logpro - pProb[iRefMap].ConstAngle[iOrient]);
pProb[iRefMap].forAngles[iOrient] += exp(logpro - pProb[iRefMap].ConstAngle[iOrient]);
/********** Getting parameters that maximize the probability ***********/
if(pProb[iRefMap].max_prob < logpro)
{
pProb[iRefMap].max_prob = logpro;
pProb[iRefMap].max_prob_cent_x = cent_x;
pProb[iRefMap].max_prob_cent_y = cent_y;
pProb[iRefMap].max_prob_orient = iOrient;
pProb[iRefMap].max_prob_conv = iConv;
}
}
/********** Getting parameters that maximize the probability ***********/
if(pProb[iRefMap].max_prob < logpro)
{
pProb[iRefMap].max_prob = logpro;
pProb[iRefMap].max_prob_cent_x = cent_x;
pProb[iRefMap].max_prob_cent_y = cent_y;
pProb[iRefMap].max_prob_orient = iOrient;
pProb[iRefMap].max_prob_conv = iConv;
}
}
}
template <int GPUAlgo, class RefT>
__device__ static inline void compareRefMapShifted(const int iRefMap, const int iOrient, const int iConv, const bioem_map& Mapconv, bioem_Probability* pProb, const bioem_param_device& param, const RefT& RefMap)
{
for (int cent_x = -param.maxDisplaceCenter; cent_x <= param.maxDisplaceCenter; cent_x=cent_x+param.GridSpaceCenter)
{
for (int cent_y = -param.maxDisplaceCenter; cent_y <= param.maxDisplaceCenter; cent_y=cent_y+param.GridSpaceCenter)
{
for (int cent_x = -param.maxDisplaceCenter; cent_x <= param.maxDisplaceCenter; cent_x=cent_x+param.GridSpaceCenter)
{
for (int cent_y = -param.maxDisplaceCenter; cent_y <= param.maxDisplaceCenter; cent_y=cent_y+param.GridSpaceCenter)
{
compareRefMap<GPUAlgo>(iRefMap, iOrient, iConv, Mapconv, pProb, param, RefMap, cent_x, cent_y);
}
}
}
}
}
#endif
include/bioem.h
View file @ 47c12f7a
......@@ -13,36 +13,44 @@
class bioem
{
public:
bioem();
virtual ~bioem();
bioem();
virtual ~bioem();
int configure(int ac, char* av[]);
int precalculate(); // Is it better to pass directly the input File names?
int dopreCalCrossCorrelation(int iRefMap, int iRefMapLocal);
int run();
int doProjections(int iMap);
int createConvolutedProjectionMap(int iOreint,int iMap, mycomplex_t* lproj,bioem_map& Mapconv,mycomplex_t* localmultFFT,myfloat_t& sumC,myfloat_t& sumsquareC);
int configure(int ac, char* av[]);
int precalculate(); // Is it better to pass directly the input File names?
int dopreCalCrossCorrelation(int iRefMap, int iRefMapLocal);
int run();
int doProjections(int iMap);
int createConvolutedProjectionMap(int iOreint,int iMap, mycomplex_t* lproj,bioem_map& Mapconv);
virtual int compareRefMaps(int iProjectionOut, int iConv, const bioem_map& conv_map, const int startMap = 0);
int createProjection(int iMap, mycomplex_t* map);
int calcross_cor(bioem_map& localmap,myfloat_t& sum,myfloat_t& sumsquare);
int compareRefMaps2(int iProjectionOut, int iConv,mycomplex_t* localmultFFT,myfloat_t sumC,myfloat_t sumsquareC);
int createProjection(int iMap, mycomplex_t* map);
int calcross_cor(bioem_map& localmap,myfloat_t& sum,myfloat_t& sumsquare);
int calProb(int iRefMap,int iOrient, int iConv,myfloat_t sumC,myfloat_t sumsquareC);
int calculateCCFFT(int iMap, int iConv, mycomplex_t* localConvFFT,mycomplex_t* localCCT,mycomplex_t* lCC);
bioem_Probability* pProb;
bioem_crossCor* crossCor;
bioem_Probability* pProb;
protected:
virtual int deviceInit();
virtual int deviceStartRun();
virtual int deviceFinishRun();
bioem_param param;
bioem_model Model;
bioem_RefMap RefMap;
bioem_param param;
bioem_model Model;
bioem_RefMap RefMap;
int nReferenceMaps; //Maps in memory at a time
int nReferenceMapsTotal; //Maps in total
int nReferenceMaps; //Maps in memory at a time
int nReferenceMapsTotal; //Maps in total
int nProjectionMaps; //Maps in memory at a time
int nProjectionMapsTotal; //Maps in total
int nProjectionMaps; //Maps in memory at a time
int nProjectionMapsTotal; //Maps in total
int FFTAlgo;
};
#endif
include/bioem_cuda_internal.h
View file @ 47c12f7a
......@@ -8,30 +8,30 @@
class bioem_cuda : public bioem
{
public:
bioem_cuda();
virtual ~bioem_cuda();
virtual int compareRefMaps(int iProjectionOut, int iConv, const bioem_map& conv_map, const int startMap = 0);
bioem_cuda();
virtual ~bioem_cuda();
virtual int compareRefMaps(int iProjectionOut, int iConv, const bioem_map& conv_map, const int startMap = 0);
protected:
virtual int deviceInit();
virtual int deviceStartRun();
virtual int deviceFinishRun();
int deviceExit();
int deviceInitialized;
cudaStream_t cudaStream;
cudaEvent_t cudaEvent[2];
bioem_RefMap_Mod* pRefMap_device_Mod;
bioem_RefMap* pRefMap_device;
bioem_Probability* pProb_device;
bioem_map* pConvMap_device[2];
int GPUAlgo; //GPU Algorithm to use, 0: parallelize over maps, 1: as 0 but work split in multiple kernels (better), 2: also parallelize over shifts (best)
int GPUAsync; //Run GPU Asynchronously, do the convolutions on the host in parallel.
int GPUWorkload; //Percentage of workload to perform on GPU. Default 100. Rest is done on processor in parallel.
int maxRef;
};
......
include/defs.h
View file @ 47c12f7a
#ifndef BIOEM_DEFS_H
#define BIOEM_DEFS_H
#ifndef BIOEM_USE_DOUBLE
typedef float myfloat_t;
typedef double mycomplex_t[2];
#define myfftw_malloc fftwf_malloc
#define myfftw_free fftwf_free
#define myfftw_destroy_plan fftwf_destroy_plan
#define myfftw_execute fftwf_execute
#define myfftw_plan_dft_2d fftwf_plan_dft_2d
#define myfftw_plan fftwf_plan
#else
typedef double myfloat_t;
#define myfftw_malloc fftw_malloc
#define myfftw_free fftw_free
#define myfftw_destroy_plan fftw_destroy_plan
#define myfftw_execute fftw_execute
#define myfftw_plan_dft_2d fftw_plan_dft_2d
#define myfftw_plan fftw_plan
#endif
typedef myfloat_t mycomplex_t[2];
#define BIOEM_FLOAT_3_PHYSICAL_SIZE 3 //Possible set to 4 for GPU
#define BIOEM_MAP_SIZE_X 224
......@@ -11,10 +27,11 @@ typedef double mycomplex_t[2];
#define BIOEM_MAX_MAPS 12000
#define MAX_REF_CTF 200
#define MAX_ORIENT 20000
#define MAX_DISPLACE 224
struct myfloat3_t
{
myfloat_t pos[BIOEM_FLOAT_3_PHYSICAL_SIZE];
myfloat_t pos[BIOEM_FLOAT_3_PHYSICAL_SIZE];
};
#ifdef BIOEM_GPUCODE
......
include/map.h
View file @ 47c12f7a
......@@ -10,30 +10,41 @@ class bioem_param;
class bioem_map
{
public:
myfloat_t points[BIOEM_MAP_SIZE_X][BIOEM_MAP_SIZE_Y];
myfloat_t points[BIOEM_MAP_SIZE_X][BIOEM_MAP_SIZE_Y];
};
class bioem_map_forFFT
{
public:
mycomplex_t cpoints[2*BIOEM_MAP_SIZE_X*2*BIOEM_MAP_SIZE_Y];
mycomplex_t cpoints[BIOEM_MAP_SIZE_X*BIOEM_MAP_SIZE_Y];
};
class bioem_convolutedMap
{
public:
bioem_map conv[MAX_REF_CTF];
myfloat_t sum_convMap[MAX_REF_CTF];
myfloat_t sumsquare_convMap[MAX_REF_CTF];
myfloat_t ForLogProbfromConv[MAX_REF_CTF];
};
class bioem_RefMap
{
public:
int readRefMaps();
const char* filemap;
int ntotRefMap;
bioem_map Ref[BIOEM_MAX_MAPS];
myfloat_t sum_RefMap[BIOEM_MAX_MAPS];
myfloat_t sumsquare_RefMap[BIOEM_MAX_MAPS];
myfloat_t ForLogProbfromRef[BIOEM_MAX_MAPS];
int readRefMaps(bioem_param& param);
int PreCalculateMapsFFT(bioem_param& param);
bioem_map_forFFT* RefMapFFT;
const char* filemap;
int ntotRefMap;
bioem_map Ref[BIOEM_MAX_MAPS];
myfloat_t sum_RefMap[BIOEM_MAX_MAPS];
myfloat_t sumsquare_RefMap[BIOEM_MAX_MAPS];
myfloat_t ForLogProbfromRef[BIOEM_MAX_MAPS];
bool dumpMap, loadMap;
__host__ __device__ inline myfloat_t get(int map, int x, int y) const {return(Ref[map].points[x][y]);}
__host__ __device__ inline const myfloat_t* getp(int map, int x, int y) const {return(&Ref[map].points[x][y]);}
};
......@@ -43,17 +54,17 @@ class bioem_RefMap_Mod
{
public:
const char* filemap;
int ntotRefMap;
myfloat_t Ref[BIOEM_MAP_SIZE_X][BIOEM_MAP_SIZE_Y][BIOEM_MAX_MAPS];
myfloat_t sum_RefMap[BIOEM_MAX_MAPS];
myfloat_t sumsquare_RefMap[BIOEM_MAX_MAPS];
myfloat_t ForLogProbfromRef[BIOEM_MAX_MAPS];
const char* filemap;
int ntotRefMap;
myfloat_t Ref[BIOEM_MAP_SIZE_X][BIOEM_MAP_SIZE_Y][BIOEM_MAX_MAPS];
myfloat_t sum_RefMap[BIOEM_MAX_MAPS];
myfloat_t sumsquare_RefMap[BIOEM_MAX_MAPS];
myfloat_t ForLogProbfromRef[BIOEM_MAX_MAPS];
__host__ __device__ inline myfloat_t get(int map, int x, int y) const {return(Ref[x][y][map]);}
bioem_RefMap_Mod() {ntotRefMap = 0;}
bioem_RefMap_Mod(const bioem_RefMap& map)
{
ntotRefMap = map.ntotRefMap;
......@@ -74,15 +85,22 @@ public:
}
};
class bioem_crossCor
{
public:
int disx[MAX_DISPLACE];
int disy[MAX_DISPLACE];
myfloat_t value[MAX_DISPLACE];
};
class bioem_Probability
{
public:
myfloat_t forAngles[MAX_ORIENT];
myfloat_t Total;
myfloat_t Constoadd;
myfloat_t ConstAngle[MAX_ORIENT];
myfloat_t max_prob;
myfloat_t forAngles[MAX_ORIENT];
myfloat_t Total;
myfloat_t Constoadd;
myfloat_t ConstAngle[MAX_ORIENT];
myfloat_t max_prob;
int max_prob_cent_x, max_prob_cent_y, max_prob_orient, max_prob_conv;
};
#endif
include/model.h
View file @ 47c12f7a
......@@ -6,22 +6,23 @@
class bioem_model
{
public:
//bioem_model();
//~bioem_model();
//bioem_model();
//~bioem_model();
int readModel();
int readModel();
bool readPDB;
bool readPDB;
myfloat_t getAminoAcidRad(char *name);
myfloat_t getAminoAcidDensity(char *name);
myfloat_t getAminoAcidRad(char *name);
myfloat_t getAminoAcidDensity(char *name);
myfloat_t NormDen;
const char* filemodel;
const char* filemodel;
int nPointsModel;
myfloat3_t PointsModel[BIOEM_MODEL_SIZE];
myfloat_t radiusPointsModel[BIOEM_MODEL_SIZE];
myfloat_t densityPointsModel[BIOEM_MODEL_SIZE];
int nPointsModel;
myfloat3_t PointsModel[BIOEM_MODEL_SIZE];
myfloat_t radiusPointsModel[BIOEM_MODEL_SIZE];
myfloat_t densityPointsModel[BIOEM_MODEL_SIZE];
};
#endif
include/param.h
View file @ 47c12f7a
......@@ -14,6 +14,7 @@ public:
int maxDisplaceCenter;
int GridSpaceCenter;
int NumberPixels;
int NtotDist;
myfloat_t Ntotpi;
myfloat_t volu;
};
......@@ -21,52 +22,52 @@ public:
class bioem_param
{
public:
bioem_param();
~bioem_param();
bioem_param();
~bioem_param();
int readParameters();
int CalculateGridsParam();
int CalculateRefCTF();
int InitializeArrays();
int readParameters();
int CalculateGridsParam();
int CalculateRefCTF();
int InitializeArrays();
bioem_param_device param_device;
bioem_map_forFFT* refCTF;
bioem_map_forFFT* refCTF;
// File names
const char* fileinput;
char logFile;
const char* fileinput;
char logFile;
// If to write Probabilities of Angles from Model
bool writeAngles;
bool writeAngles;
// Pixel size && BIOEM_MAP_SIZE_X should be defined here too
//int NumberPixels; //in device class
myfloat_t* pixelSize;
//int NumberPixels; //in device class
myfloat_t pixelSize;
// Grid Points in Euler angles, assuming uniform sampling d_alpha=d_gamma (in 2pi) & cos(beta)=-1,1
int angleGridPointsAlpha;
int angleGridPointsBeta;
int angleGridPointsAlpha;
int angleGridPointsBeta;
// Grids in center assuming equidistance from 0,0
//int maxDisplaceCenter; //in device class
int numberGridPointsDisplaceCenter;
//int maxDisplaceCenter; //in device class
int numberGridPointsDisplaceCenter;
//int GridSpaceCenter; //in device class
// Grid sampling for the convolution kernel
//ENVELOPE
myfloat_t* startGridEnvelop;
int numberGridPointsEnvelop;
myfloat_t* gridEnvelop;
//CTF=Amp*cos(phase*x)-sqrt(1-Amp**2)*sin(phase*x)
myfloat_t* startGridCTF_phase;
int numberGridPointsCTF_phase;
myfloat_t* gridCTF_phase;
myfloat_t* startGridCTF_amp;
int numberGridPointsCTF_amp;
myfloat_t* gridCTF_amp;
// Others
//myfloat_t volu;//in device class
myfloat3_t angles[MAX_ORIENT];
myfloat3_t CtfParam[MAX_REF_CTF] ;
int nTotGridAngles;
int nTotCTFs;
//myfloat_t Ntotpi;//in device class
//ENVELOPE
myfloat_t startGridEnvelop;
int numberGridPointsEnvelop;
myfloat_t gridEnvelop;
//CTF=Amp*cos(phase*x)-sqrt(1-Amp**2)*sin(phase*x)
myfloat_t startGridCTF_phase;
int numberGridPointsCTF_phase;
myfloat_t gridCTF_phase;
myfloat_t startGridCTF_amp;
int numberGridPointsCTF_amp;
myfloat_t gridCTF_amp;
// Others
//myfloat_t volu;//in device class
myfloat3_t angles[MAX_ORIENT];
myfloat3_t CtfParam[MAX_REF_CTF] ;
int nTotGridAngles;
int nTotCTFs;
//myfloat_t Ntotpi;//in device class
};
#endif
main.cpp
View file @ 47c12f7a
......@@ -21,41 +21,41 @@
int main(int argc, char* argv[])
{
/**************************************************************************************/
/********************************* Main BioEM code **********************************/
/************************************************************************************/
/**************************************************************************************/
/********************************* Main BioEM code **********************************/
/************************************************************************************/
#pragma omp parallel
{
_MM_SET_DENORMALS_ZERO_MODE(_MM_DENORMALS_ZERO_ON); //Flush denormals to zero in all OpenMP threads
}
HighResTimer timer;
HighResTimer timer;
bioem* bio;
bioem* bio;
#ifdef WITH_CUDA
if (getenv("GPU") && atoi(getenv("GPU")))
{
if (getenv("GPU") && atoi(getenv("GPU")))
{
bio = bioem_cuda_create();
}
else
}
else
#endif
{
bio = new bioem;
{
bio = new bioem;
}
/************ Configuration and Pre-calculating necessary objects *****************/
printf("Configuring\n");
bio->configure(argc,argv);
/************ Configuration and Pre-calculating necessary objects *****************/
printf("Configuring\n");
bio->configure(argc,argv);
/******************************* Run BioEM routine ******************************/
printf("Running\n");
timer.Start();
bio->run();
timer.Stop();
/******************************* Run BioEM routine ******************************/
printf("Running\n");
timer.Start();
bio->run();
timer.Stop();
/************************************ End **********************************/
printf ("The code ran for %f seconds.\n", timer.GetElapsedTime());
delete bio;
/************************************ End **********************************/
printf ("The code ran for %f seconds.\n", timer.GetElapsedTime());
delete bio;
return(0);
return(0);
}
map.cpp
View file @ 47c12f7a
......@@ -3,16 +3,20 @@
#include <stdio.h>
#include <stdlib.h>
#include <cstring>
#include <math.h>
#include <fftw3.h>
#include "map.h"
#include "param.h"
using namespace std;
int bioem_RefMap::readRefMaps()
int bioem_RefMap::readRefMaps(bioem_param& param)
{
/**************************************************************************************/
/***********************Reading reference Particle Maps************************/
/**************************************************************************************/
/**************************************************************************************/
/***********************Reading reference Particle Maps************************/
/**************************************************************************************/
if (loadMap)
{
FILE* fp = fopen("maps.dump", "rb");
......@@ -29,13 +33,15 @@ int bioem_RefMap::readRefMaps()
}
fread(&Ref[0], sizeof(Ref[0]), ntotRefMap, fp);
fclose(fp);
cout << "Particle Maps read from Map Dump \nTotal Number of particles: " << ntotRefMap ;
cout << "\n+++++++++++++++++++++++++++++++++++++++++ \n";
cout << "\n+++++++++++++++++++++++++++++++++++++++++ \n";
}
else
{
int nummap=-1;
int lasti=0;
int lastj=0;
ifstream input(filemap);
if (!input.good())
{
......@@ -65,12 +71,17 @@ int bioem_RefMap::readRefMaps()
cout << "BIOEM_MAX_MAPS too small\n";
exit(1);
}
if(lasti!=param.param_device.NumberPixels && lastj!=param.param_device.NumberPixels && nummap>0)
{
cout << "PROBLEM INCONSISTENT NUMBER OF PIXELS IN MAPS AND INPUTFILE ( " << param.param_device.NumberPixels << ", i " << lasti << ", j " << lastj << ")" << "\n";
exit(1);
}
}
else
{
int i,j;
myfloat_t z;
char tmpVals[36] = {' '};
strncpy (tmpVals,line,8);
......@@ -85,6 +96,8 @@ int bioem_RefMap::readRefMaps()
if(i>0 && i-1 <BIOEM_MAP_SIZE_X&& j>0 && j-1<BIOEM_MAP_SIZE_Y && nummap < BIOEM_MAX_MAPS)
{
Ref[nummap].points[i-1][j-1]=z;
lasti=i;
lastj=j;
}
else
{
......@@ -97,7 +110,7 @@ int bioem_RefMap::readRefMaps()
ntotRefMap=nummap+1;
cout << "Particle Maps read from Standard File \nTotal Number of particles: " << ntotRefMap ;
cout << "\n+++++++++++++++++++++++++++++++++++++++++ \n";
if (dumpMap)
{
FILE* fp = fopen("maps.dump", "w+b");
......@@ -112,5 +125,60 @@ int bioem_RefMap::readRefMaps()
}
}
return(0);
return(0);
}
int bioem_RefMap::PreCalculateMapsFFT(bioem_param& param)
{
/**************************************************************************************/
/********** Routine that pre-calculates Kernels for Convolution **********************/
/************************************************************************************/
myfftw_plan plan;
mycomplex_t* localMap;