bool mscore_kgpu::add_mi(mspectrum &_s)
{
if (!mscore::add_mi(_s))
return false;
if (&_s == m_vSpectraToScore.back()) { // last in list, transfer
// all to device memory in
// one go
CUDA_TIMER_START(mi_time);
size_t a, n_pairs=0;
for (a=0;a < m_vSpectraToScore.size();a++) {
n_pairs += m_vSpectraToScore[a]->m_vMI.size();
}
pinned_host_vector_float_t fM(n_pairs);
pinned_host_vector_float_t fI(n_pairs);
n_pairs = 0;
for (a=0;a < m_vSpectraToScore.size();a++) {
for (size_t n=0;n<m_vSpectraToScore[a]->m_vMI.size();n++) {
fI[n_pairs] = m_vSpectraToScore[a]->m_vMI[n].m_fI;
fM[n_pairs++] = m_vSpectraToScore[a]->m_vMI[n].m_fM;
}
}
// now copy to memory
thrust::device_vector<float> dfM;
mscore_kgpu_thrust_host_to_device_copy_float(fM,dfM);
thrust::device_vector<float> dfI;
mscore_kgpu_thrust_host_to_device_copy_float(fI,dfI);
// and actually do the add_mi logic
n_pairs = 0;
for (a=0;a < m_vSpectraToScore.size();a++) {
mspectrum &_s = *m_vSpectraToScore[a];
vmiTypeGPU vTypeGPU;
vTypeGPU.init(_s.m_vMI.size());
if (_s.m_vMI.size() != 0)
{
// Screen peeks on upper end.
int iWindowCount = 10;
int endMassMax = (int)(
((_s.m_dMH + (_s.m_fZ - 1) * m_seqUtil.m_dProton) / _s.m_fZ)
*
2.0 + 0.5)
+
iWindowCount;
// now pass off to CUDA implementation
mscore_kgpu_thrust_score(
dfM.begin()+n_pairs,
dfI.begin()+n_pairs,_s.m_vMI.size(),
m_dIsotopeCorrection,
iWindowCount,
endMassMax,
m_maxEnd,
vTypeGPU); // results come back in vTypeGPU
}
m_vSpectraGPU.push_back(vTypeGPU);
n_pairs += _s.m_vMI.size();
}
CUDA_TIMER_STOP(mi_time)
}
int main() {
typedef Kokkos::DefaultNode::DefaultNodeType Node;
typedef KokkosExamples::DummySparseKernel<Node> SparseOps;
typedef Kokkos::CrsGraph < int,Node,SparseOps> Graph;
typedef Kokkos::CrsMatrix<double,int,Node,SparseOps> DoubleMat;
typedef Kokkos::CrsMatrix< float,int,Node,SparseOps> FloatMat;
typedef Kokkos::MultiVector<double,Node> DoubleVec;
typedef Kokkos::MultiVector<float,Node> FloatVec;
std::cout << "Note, this class doesn't actually do anything. We are only testing that it compiles." << std::endl;
// get a pointer to the default node
Teuchos::RCP<Node> node = Kokkos::DefaultNode::getDefaultNode();
// create the graph G
const size_t numRows = 5;
Graph G(numRows,node);
// create a double-valued matrix dM using the graph G
DoubleMat dM(G);
// create a double-valued sparse kernel using the rebind functionality
SparseOps::rebind<double>::other doubleKernel(node);
// initialize it with G and dM
doubleKernel.initializeStructure(G);
doubleKernel.initializeValues(dM);
// create double-valued vectors and initialize them
DoubleVec dx(node), dy(node);
// test the sparse kernel operator interfaces
doubleKernel.multiply( Teuchos::NO_TRANS, 1.0, dx, dy);
doubleKernel.multiply( Teuchos::NO_TRANS, 1.0, dx, 1.0, dy);
doubleKernel.solve( Teuchos::NO_TRANS, Teuchos::UPPER_TRI, Teuchos::UNIT_DIAG, dy, dx);
// create a float-valued matrix fM using the graph G
FloatMat fM(G);
// create a double-valued sparse kernel using the rebind functionality
SparseOps::rebind<float>::other floatKernel(node);
// initialize it with G and fM
floatKernel.initializeStructure(G);
floatKernel.initializeValues(fM);
// create float-valued vectors and initialize them
FloatVec fx(node), fy(node);
// test the sparse kernel operator interfaces
floatKernel.multiply( Teuchos::NO_TRANS, 1.0f, fx, fy);
floatKernel.multiply( Teuchos::NO_TRANS, 1.0f, fx, 1.0f, fy);
floatKernel.solve( Teuchos::NO_TRANS, Teuchos::UPPER_TRI, Teuchos::UNIT_DIAG, fy, fx);
std::cout << "End Result: TEST PASSED" << std::endl;
return 0;
}
cMatrCorresp * cModeleAnalytiqueComp::GetMatr(int aPas,bool PointUnique)
{
int aDz = mEtape->DeZoomTer();
mSzGl = mAppli.SzOfResol(aDz);
mSz = (mSzGl+Pt2di(aPas-1,aPas-1))/aPas;
cMatrCorresp * pMatr = new cMatrCorresp (mAppli,mSz,mNbPx);
Fonc_Num aFPds = mAppli.FoncMasqOfResol(aDz);
if (mAppli.OneDefCorAllPxDefCor().Val())
aFPds = aFPds && mEtape->FileMasqOfNoDef().in();
if (PointUnique)
{
aFPds = aFPds * ((FX%aPas)==(aPas/2)) * ((FY%aPas)==(aPas/2));
}
// std::cout << "AAAAAAAAAa\n";
if (mModele.FiltreByCorrel().Val())
{
Im1D_REAL8 aLut(256);
double aS = mModele.SeuilFiltreCorrel().Val();
double anExp = mModele.ExposantPondereCorrel().Val();
Fonc_Num FC = FX/128.0-1;
ELISE_COPY
(
aLut.all_pts(),
(mModele.UseFCBySeuil().Val()) ?
(FC>=aS) :
pow(Max(0.0,FC-aS),anExp),
aLut.out()
);
aFPds = aFPds * aLut.in()[mEtape->LastFileCorrelOK().in()];
/*
double aS = mModele.SeuilFiltreCorrel().Val();
Fonc_Num FC = mEtape->LastFileCorrelOK().in()/128.0-1;
if (mModele.UseFCBySeuil().Val())
aFPds = aFPds * (FC>=aS);
else
aFPds = aFPds * Max(0.0,(FC-aS));
*/
}
Symb_FNum fM(aFPds);
/*
{
double aSom;
ELISE_COPY
(
rectangle(Pt2di(0,0),mSzGl),
fM,
sigma(aSom)
);
std::cout << "SssDSOM = " << aSom << "\n";
getchar();
}
*/
Fonc_Num aFonc = Virgule(fM,fM*FX,fM*FY);
for (int aK=0; aK<mNbPx ; aK++)
{
aFonc = Virgule(aFonc,fM*ImPx(aK));
}
ELISE_COPY
(
rectangle(Pt2di(0,0),mSzGl),
aFonc,
pMatr->StdHisto().chc(Virgule(FX,FY)/aPas)
);
pMatr->Normalize(mModele,mGeoTer,mPDV1,mGeom1,mPDV2,mGeom2);
return pMatr;
}
