Skip to content

GitLab

Commit b9e18ed6 authored by wangkx1's avatar wangkx1
Browse files

init

parents
Hide whitespace changes
Inline Side-by-side
Showing with 5761 additions and 0 deletions
cuda/StackResampler.cpp 0 → 100644
View file @ b9e18ed6
/*! \file StackResampler.cu
\brief Contains definition of CUDA implementation of a class for spline resampling of irregularly sampled columns in the z-direction
\author Jesper Andersson
\version 1.0b, March, 2016.
*/
//
// StackResampler.cu
//
// Jesper Andersson, FMRIB Image Analysis Group
//
// Copyright (C) 2016 University of Oxford
//
// Because of a bug in cuda_fp16.hpp, that gets included by hipblas.h, it has to
// be included before any include files that set up anything related to the std-lib.
// If not, there will be an ambiguity in cuda_fp16.hpp about wether to use the
// old-style C isinf or the new (since C++11) std::isinf.
#include "hipblas.h"
#include <cstdlib>
#include <string>
#include <vector>
#include <cmath>
#include <hip/hip_runtime.h>
#include <thrust/system_error.h>
#include <thrust/host_vector.h>
#include <thrust/device_vector.h>
#include <thrust/device_ptr.h>
#include <thrust/transform.h>
#include <thrust/fill.h>
#pragma push
#pragma diag_suppress = code_is_unreachable // Supress warnings from armawrap
#include "armawrap/newmat.h"
#include "newimage/newimage.h"
#include "EddyHelperClasses.h"
#include "CudaVolume.h"
#pragma pop
#include "EddyCudaHelperFunctions.h"
#include "StackResampler.h"
#include "EddyKernels.h"
#include "EddyMatrixKernels.h"
#include "EddyFunctors.h"
namespace EDDY {
const dim3 StackResampler::_threads_per_block_WtW_StS = dim3(16,16);
StackResampler::StackResampler(const EDDY::CudaVolume& stack,
const EDDY::CudaVolume& zcoord,
const EDDY::CudaVolume& pred,
const EDDY::CudaVolume& mask,
double lambda) EddyTry : _resvol(stack,false), _mask(stack,false)
{
// static unsigned int cnt = 0;
unsigned int matsz = stack.Size(2);
unsigned int nmat = stack.Size(0); // Number of matrices per xz-plane
// Allocate memory for storing matrices/vectors on the GPU
thrust::device_vector<float> StS_matrix(sqr(matsz));
thrust::device_vector<float> empty_StS_matrix(sqr(matsz));
thrust::device_vector<float> gpu_W_matrix(sqr(matsz));
thrust::device_vector<float> gpu_sorted_zcoord(zcoord.Size());
thrust::device_vector<float> gpu_Wir_matrices(nmat*sqr(matsz));
thrust::device_vector<float> gpu_weights(nmat*matsz);
thrust::device_vector<float> gpu_diagw_p_vectors(nmat*matsz);
thrust::device_vector<float> gpu_diagw_W_matrices(nmat*sqr(matsz));
thrust::device_vector<float> gpu_Wirty_vectors(nmat*matsz);
thrust::device_vector<float> gpu_dWtdp_vectors(nmat*matsz);
thrust::device_vector<float> gpu_WtW_matrices(nmat*sqr(matsz));
thrust::device_vector<float> gpu_dWtdW_matrices(nmat*sqr(matsz));
thrust::device_vector<float> gpu_sum_vectors(nmat*matsz);
thrust::device_vector<float> gpu_sum_matrices(nmat*sqr(matsz));
thrust::device_vector<float> gpu_c_hat_vectors(nmat*matsz);
thrust::device_vector<float> gpu_y_hat_vectors(nmat*matsz);
// EDDY::CudaVolume weights = pred;
// weights = 0.0;
// Get regularisation matrix once and for all
//cout << "Making reg matrix" << endl; cout.flush();
get_StS(matsz,lambda,StS_matrix);
get_StS(matsz,0.0,empty_StS_matrix);
// Get weight matrix for regular grid once and for all
//cout << "Making W" << endl; cout.flush();
get_regular_W(matsz,gpu_W_matrix);
// Populate mask
//cout << "Making mask" << endl; cout.flush();
make_mask(mask,zcoord,true,_mask);
// Get a volume of sorted z-ccords
//cout << "sorting zcoords" << endl; cout.flush();
sort_zcoords(zcoord,true,gpu_sorted_zcoord);
// Interpolate along z-columns
for (unsigned int j=0; j<stack.Size(1); j++) { // Loop over all xz-planes
// Make all Wir matrices for this xz-plane
//cout << "Making Wir" << endl; cout.flush();
make_Wir_matrices(zcoord,j,true,gpu_Wir_matrices);
// Make vectors of weights for predictions
//cout << "Making weights" << endl; cout.flush();
make_weight_vectors(gpu_sorted_zcoord,stack.Size(0),stack.Size(2),j,true,gpu_weights);
// insert_weights(gpu_weights,j,true,weights);
// Make diag{w}p vectors, where p are the predictions
//cout << "Making diag{w}p" << endl; cout.flush();
make_diagw_p_vectors(pred,gpu_weights,j,true,gpu_diagw_p_vectors);
// Make diag{w}W matrices, where W is the spline design matrix for regular sampling
//cout << "Making diag{w}W" << endl; cout.flush();
make_diagw_W_matrices(gpu_weights,gpu_W_matrix,matsz,nmat,true,gpu_diagw_W_matrices);
// Calculate all Wir'*y vectors for this plane
//cout << "Making Wirty" << endl; cout.flush();
make_Wir_t_y_vectors(stack,gpu_Wir_matrices,j,true,gpu_Wirty_vectors);
// Calculate all (diag{w}W)'*diag{w}p vectors for this plane
//cout << "Making diag{w}W'diag{w}p" << endl; cout.flush();
make_dwWt_dwp_vectors(gpu_diagw_W_matrices,gpu_diagw_p_vectors,matsz,nmat,true,gpu_dWtdp_vectors);
// Calculate all Wir'*Wir matrices
//cout << "Making Wir'Wir" << endl; cout.flush();
make_WtW_StS_matrices(gpu_Wir_matrices,matsz,nmat,StS_matrix,true,gpu_WtW_matrices);
// Calculate all (diag{w}W)'*(diag{w}W) matrices
//cout << "Making diag{w}W'diag{w}W" << endl; cout.flush();
make_WtW_StS_matrices(gpu_diagw_W_matrices,matsz,nmat,empty_StS_matrix,true,gpu_dWtdW_matrices);
// Add (pairwise) all Wir'*y and (diag{w}W)'*diag{w}p vectors
//cout << "Making sum vectors" << endl; cout.flush();
thrust::transform(gpu_Wirty_vectors.begin(),gpu_Wirty_vectors.end(),gpu_dWtdp_vectors.begin(),gpu_sum_vectors.begin(),thrust::plus<float>());
// Add (pairwise) all Wir'*Wir and (diag{w}W)'*(diag{w}W)
//cout << "Making sum matrices" << endl; cout.flush();
thrust::transform(gpu_WtW_matrices.begin(),gpu_WtW_matrices.end(),gpu_dWtdW_matrices.begin(),gpu_sum_matrices.begin(),thrust::plus<float>());
// Solve for spline coefficients
//cout << "Making c_hat" << endl; cout.flush();
solve_for_c_hat(gpu_sum_matrices,gpu_sum_vectors,matsz,nmat,true,gpu_c_hat_vectors);
// Multiply by W to solve for y_hat (interpolate in z-direction)
//cout << "Making y_hat" << endl; cout.flush();
make_y_hat_vectors(gpu_W_matrix,gpu_c_hat_vectors,matsz,nmat,true,gpu_y_hat_vectors);
// Transfer solution vectors to result volume
//cout << "Transferring y_hat" << endl; cout.flush();
transfer_y_hat_to_volume(gpu_y_hat_vectors,j,true,_resvol);
/*
if (cnt==19 && j==30) {
write_debug_info_for_pred_resampling(54,j,"debug_info.txt",zcoord,stack,pred,gpu_sorted_zcoord,
gpu_W_matrix,gpu_Wir_matrices,gpu_weights,gpu_diagw_p_vectors,
gpu_diagw_W_matrices,gpu_Wirty_vectors,gpu_dWtdp_vectors,
gpu_WtW_matrices,gpu_dWtdW_matrices,gpu_sum_vectors,
gpu_sum_matrices,gpu_c_hat_vectors,gpu_y_hat_vectors);
}
*/
}
// char fname[256]; sprintf(fname,"weights_%02d",cnt);
// weights.Write(std::string(fname));
// cnt++;
} EddyCatch
StackResampler::StackResampler(const EDDY::CudaVolume& stack,
const EDDY::CudaVolume& zcoord,
const EDDY::CudaVolume& mask,
NEWIMAGE::interpolation interp,
double lambda) EddyTry : _resvol(stack,false), _mask(stack,false)
{
// Populate mask
make_mask(mask,zcoord,true,_mask);
if (interp == NEWIMAGE::spline) {
unsigned int matsz = stack.Size(2);
unsigned int nmat = stack.Size(0); // Number of matrices per xz-plane
// Allocate memory for storing matrices/vectors on the GPU
thrust::device_vector<float> gpu_StS_matrix(sqr(matsz));
thrust::device_vector<float> gpu_W_matrix(sqr(matsz));
thrust::device_vector<float> gpu_Wir_matrices(nmat*sqr(matsz));
thrust::device_vector<float> gpu_WtW_StS_matrices(nmat*sqr(matsz));
thrust::device_vector<float> gpu_Wirty_vectors(nmat*matsz);
thrust::device_vector<float> gpu_c_hat_vectors(nmat*matsz);
thrust::device_vector<float> gpu_y_hat_vectors(nmat*matsz);
// Get regularisation matrix once and for all
get_StS(matsz,lambda,gpu_StS_matrix);
// Get spline matrix for regular grid once and for all
get_regular_W(matsz,gpu_W_matrix);
// Interpolate along z-columns
for (unsigned int j=0; j<stack.Size(1); j++) { // Loop over all xz-planes
// Make all Wir matrices for this xz-plane
make_Wir_matrices(zcoord,j,true,gpu_Wir_matrices);
// Pre-multiply Wir by transpose of self and add lambda*StS.
make_WtW_StS_matrices(gpu_Wir_matrices,matsz,nmat,gpu_StS_matrix,true,gpu_WtW_StS_matrices);
// Multiply Wir transposed with y (observed intensities).
make_Wir_t_y_vectors(stack,gpu_Wir_matrices,j,true,gpu_Wirty_vectors);
// Solve for spline coefficients
solve_for_c_hat(gpu_WtW_StS_matrices,gpu_Wirty_vectors,matsz,nmat,true,gpu_c_hat_vectors);
// Multiply by W to solve for y_hat (interpolate in z-direction)
make_y_hat_vectors(gpu_W_matrix,gpu_c_hat_vectors,matsz,nmat,true,gpu_y_hat_vectors);
// Transfer solution vectors to result volume
transfer_y_hat_to_volume(gpu_y_hat_vectors,j,true,_resvol);
}
}
else if (interp == NEWIMAGE::trilinear) {
thrust::device_vector<float> gpu_sorted_zcoord(zcoord.Size());
thrust::device_vector<float> gpu_sorted_intensities(zcoord.Size());
thrust::device_vector<float> gpu_interpolated_columns(zcoord.Size(),0.0);
// Sort the z-ccordinates and the intensities with z-coord as key
sort_zcoords_and_intensities(zcoord,stack,true,gpu_sorted_zcoord,gpu_sorted_intensities);
// Do linear interpolation along z-columns
linear_interpolate_columns(gpu_sorted_zcoord,gpu_sorted_intensities,zcoord.Size(0),zcoord.Size(1),zcoord.Size(2),true,gpu_interpolated_columns);
// Transfer interpolated vectors to results volume
transfer_interpolated_columns_to_volume(gpu_interpolated_columns,true,_resvol);
}
else throw EddyException("StackResampler::StackResampler: Invalid interpolation method");
} EddyCatch
void StackResampler::get_StS(unsigned int sz, double lambda, thrust::device_vector<float>& StS) const EddyTry
{
float six = lambda * 6.0; float minusfour = - lambda * 4.0; float one = lambda;
thrust::host_vector<float> hStS(StS.size(),0.0);
hStS[rfindx(0,0,sz)] = six; hStS[rfindx(0,1,sz)] = minusfour; hStS[rfindx(0,2,sz)] = one; hStS[rfindx(0,(sz-2),sz)] = one; hStS[rfindx(0,(sz-1),sz)] = minusfour;
hStS[rfindx(1,0,sz)] = minusfour; hStS[rfindx(1,1,sz)] = six; hStS[rfindx(1,2,sz)] = minusfour; hStS[rfindx(1,3,sz)] = one; hStS[rfindx(1,(sz-1),sz)] = one;
for ( unsigned int i=2; i<(sz-2); i++) {
hStS[rfindx(i,i-2,sz)] = one;
hStS[rfindx(i,i-1,sz)] = minusfour;
hStS[rfindx(i,i,sz)] = six;
hStS[rfindx(i,i+1,sz)] = minusfour;
hStS[rfindx(i,i+2,sz)] = one;
}
hStS[rfindx((sz-2),0,sz)] = one; hStS[rfindx((sz-2),(sz-4),sz)] = one; hStS[rfindx((sz-2),(sz-3),sz)] = minusfour; hStS[rfindx((sz-2),(sz-2),sz)] = six; hStS[rfindx((sz-2),(sz-1),sz)] = minusfour;
hStS[rfindx((sz-1),0,sz)] = minusfour; hStS[rfindx((sz-1),1,sz)] = one; hStS[rfindx((sz-1),(sz-3),sz)] = one; hStS[rfindx((sz-1),(sz-2),sz)] = minusfour; hStS[rfindx((sz-1),(sz-1),sz)] = six;
StS = hStS;
return;
} EddyCatch
void StackResampler::get_regular_W(unsigned int sz, thrust::device_vector<float>& W) const EddyTry
{
thrust::host_vector<float> hW(W.size(),0.0);
hW[rfindx(0,0,sz)] = 5.0/6.0; hW[rfindx(0,1,sz)] = 1.0/6.0;
for ( unsigned int i=1; i<(sz-1); i++) {
hW[rfindx(i,i-1,sz)] = 1.0/6.0;
hW[rfindx(i,i,sz)] = 4.0/6.0;
hW[rfindx(i,i+1,sz)] = 1.0/6.0;
}
hW[rfindx((sz-1),(sz-2),sz)] = 1.0/6.0; hW[rfindx((sz-1),(sz-1),sz)] = 5.0/6.0;
W = hW;
return;
} EddyCatch
void StackResampler::make_mask(const EDDY::CudaVolume& inmask,
const EDDY::CudaVolume& zcoord,
bool sync,
EDDY::CudaVolume& omask) EddyTry
{
omask = 0.0;
int nblocks = inmask.Size(1);
int tpb = inmask.Size(0);
EddyKernels::make_mask_from_stack<<<nblocks,tpb>>>(inmask.GetPtr(),zcoord.GetPtr(),inmask.Size(0),
inmask.Size(1),inmask.Size(2),omask.GetPtr());
if (sync) EddyCudaHelperFunctions::CudaSync("EddyKernels::make_mask_from_stack");
} EddyCatch
void StackResampler::sort_zcoords(const EDDY::CudaVolume& zcoord,
bool sync,
thrust::device_vector<float>& szcoord) const EddyTry
{
// First transfer all columns and check if they are sorted
thrust::device_vector<unsigned int> ns_flags(zcoord.Size(0)*zcoord.Size(1),0); // Non-zero if a column is "not sorted"
EddyKernels::TransferAndCheckSorting<<<zcoord.Size(1),zcoord.Size(0)>>>(zcoord.GetPtr(),zcoord.Size(0),
zcoord.Size(1),zcoord.Size(2),
thrust::raw_pointer_cast(szcoord.data()),
thrust::raw_pointer_cast(ns_flags.data()));
EddyCudaHelperFunctions::CudaSync("EddyKernels::TransferAndCheckSorting");
unsigned int nnsort = thrust::reduce(ns_flags.begin(),ns_flags.end(),0,EDDY::Sum<unsigned int,unsigned int>());
if (nnsort) { // If there are columns that are not sorted
thrust::host_vector<unsigned int> host_ns_flags = ns_flags;
thrust::host_vector<unsigned int> host_nsort_indx(nnsort,0);
for (unsigned int i=0, n=0; i<zcoord.Size(0)*zcoord.Size(1); i++) {
if (host_ns_flags[i]) { host_nsort_indx[n] = i; n++; }
}
thrust::device_vector<unsigned int> device_nsort_indx = host_nsort_indx;
int nb = (nnsort / 32) + 1;
EddyKernels::SortVectors<<<nb,32>>>(thrust::raw_pointer_cast(device_nsort_indx.data()),nnsort,
zcoord.Size(2),thrust::raw_pointer_cast(szcoord.data()),NULL);
if (sync) EddyCudaHelperFunctions::CudaSync("EddyKernels::SortVectors");
/*
// Check that vectors really are sorted. Only for use during testing.
thrust::host_vector<float> hvec = szcoord;
for (unsigned int vi=0; vi<zcoord.Size(1)*zcoord.Size(0); vi++) {
unsigned int offs = vi*zcoord.Size(2);
for (unsigned int i=1; i<zcoord.Size(2); i++) {
if (hvec[offs+i] < hvec[offs+i-1]) {
cout << "Vector " << vi << " was not sucessfully sorted" << endl;
cout.flush();
for (unsigned int j=0; j<zcoord.Size(2); j++) {
cout << "hvec[" << j << "] = " << hvec[offs+j] << endl;
}
exit(0);
}
}
}
*/
}
} EddyCatch
void StackResampler::sort_zcoords_and_intensities(const EDDY::CudaVolume& zcoord,
const EDDY::CudaVolume& data,
bool sync,
thrust::device_vector<float>& szcoord,
thrust::device_vector<float>& sdata) const EddyTry
{
// First transfer all zcoord columns and check if they are sorted
thrust::device_vector<unsigned int> ns_flags(zcoord.Size(0)*zcoord.Size(1),0); // Non-zero if a column is "not sorted"
EddyKernels::TransferAndCheckSorting<<<zcoord.Size(1),zcoord.Size(0)>>>(zcoord.GetPtr(),zcoord.Size(0),
zcoord.Size(1),zcoord.Size(2),
thrust::raw_pointer_cast(szcoord.data()),
thrust::raw_pointer_cast(ns_flags.data()));
EddyCudaHelperFunctions::CudaSync("EddyKernels::TransferAndCheckSorting");
unsigned int nnsort = thrust::reduce(ns_flags.begin(),ns_flags.end(),0,EDDY::Sum<unsigned int,unsigned int>());
// Transfer all intensity values as well
EddyKernels::TransferVolumeToVectors<<<data.Size(1),data.Size(0)>>>(data.GetPtr(),data.Size(0),
data.Size(1),data.Size(2),
thrust::raw_pointer_cast(sdata.data()));
if (nnsort) { // If there are columns that are not sorted
thrust::host_vector<unsigned int> host_ns_flags = ns_flags;
thrust::host_vector<unsigned int> host_nsort_indx(nnsort,0);
for (unsigned int i=0, n=0; i<zcoord.Size(0)*zcoord.Size(1); i++) {
if (host_ns_flags[i]) { host_nsort_indx[n] = i; n++; }
}
thrust::device_vector<unsigned int> device_nsort_indx = host_nsort_indx;
int nb = (nnsort / 32) + 1;
EddyKernels::SortVectors<<<nb,32>>>(thrust::raw_pointer_cast(device_nsort_indx.data()),nnsort,
zcoord.Size(2),thrust::raw_pointer_cast(szcoord.data()),
thrust::raw_pointer_cast(sdata.data()));
if (sync) EddyCudaHelperFunctions::CudaSync("EddyKernels::SortVectors");
}
} EddyCatch
void StackResampler::linear_interpolate_columns(const thrust::device_vector<float>& zcoord,
const thrust::device_vector<float>& val,
unsigned int xsz,
unsigned int ysz,
unsigned int zsz,
bool sync,
thrust::device_vector<float>& ival) const EddyTry
{
EddyKernels::LinearInterpolate<<<ysz,xsz>>>(thrust::raw_pointer_cast(zcoord.data()),
thrust::raw_pointer_cast(val.data()),zsz,
thrust::raw_pointer_cast(ival.data()));
if (sync) EddyCudaHelperFunctions::CudaSync("EddyKernels::LinearInterpolate");
} EddyCatch
void StackResampler::transfer_interpolated_columns_to_volume(const thrust::device_vector<float>& zcols,
bool sync,
EDDY::CudaVolume& vol) EddyTry
{
EddyKernels::TransferColumnsToVolume<<<vol.Size(1),vol.Size(0)>>>(thrust::raw_pointer_cast(zcols.data()),
vol.Size(0),vol.Size(1),
vol.Size(2),vol.GetPtr());
if (sync) EddyCudaHelperFunctions::CudaSync("EddyKernels::TransferColumnsToVolume");
} EddyCatch
void StackResampler::make_weight_vectors(const thrust::device_vector<float>& zcoord,
unsigned int xsz,
unsigned int zsz,
unsigned int xzp,
bool sync,
thrust::device_vector<float>& weights) const EddyTry
{
EddyKernels::MakeWeights<<<xsz,zsz>>>(thrust::raw_pointer_cast(zcoord.data()),xsz,
zsz,xzp,thrust::raw_pointer_cast(weights.data()));
if (sync) EddyCudaHelperFunctions::CudaSync("EddyKernels::MakeWeights");
} EddyCatch
void StackResampler::insert_weights(const thrust::device_vector<float>& wvec,
unsigned int j,
bool sync,
EDDY::CudaVolume& wvol) const EddyTry
{
EddyKernels::InsertWeights<<<wvol.Size(0),wvol.Size(2)>>>(thrust::raw_pointer_cast(wvec.data()),j,wvol.Size(0),
wvol.Size(1),wvol.Size(2),wvol.GetPtr());
if (sync) EddyCudaHelperFunctions::CudaSync("EddyKernels::InsertWeights");
} EddyCatch
void StackResampler::make_diagw_p_vectors(const EDDY::CudaVolume& pred,
const thrust::device_vector<float>& wgts,
unsigned int xzp,
bool sync,
thrust::device_vector<float>& wp) const EddyTry
{
EddyKernels::MakeDiagwpVecs<<<pred.Size(0),pred.Size(2)>>>(pred.GetPtr(),thrust::raw_pointer_cast(wgts.data()),
pred.Size(0),pred.Size(1),pred.Size(2),xzp,
thrust::raw_pointer_cast(wp.data()));
if (sync) EddyCudaHelperFunctions::CudaSync("EddyKernels::MakeDiagwpVecs");
} EddyCatch
void StackResampler::make_diagw_W_matrices(const thrust::device_vector<float>& wgts,
const thrust::device_vector<float>& W,
unsigned int matsz,
unsigned int nmat,
bool sync,
thrust::device_vector<float>& diagwW) const EddyTry
{
EddyMatrixKernels::DiagwA<<<nmat,matsz>>>(thrust::raw_pointer_cast(wgts.data()),
thrust::raw_pointer_cast(W.data()),
matsz,matsz,nmat,
thrust::raw_pointer_cast(diagwW.data()));
if (sync) EddyCudaHelperFunctions::CudaSync("EddyMatrixKernels::DiagwA");
} EddyCatch
void StackResampler::make_dwWt_dwp_vectors(const thrust::device_vector<float>& dW,
const thrust::device_vector<float>& dp,
unsigned int matsz,
unsigned int nmat,
bool sync,
thrust::device_vector<float>& dWtdp) const EddyTry
{
EddyMatrixKernels::Atb<<<nmat,matsz>>>(thrust::raw_pointer_cast(dW.data()),
thrust::raw_pointer_cast(dp.data()),
matsz,matsz,nmat,nmat,
thrust::raw_pointer_cast(dWtdp.data()));
if (sync) EddyCudaHelperFunctions::CudaSync("EddyMatrixKernels::Atb");
return;
} EddyCatch
void StackResampler::make_Wir_matrices(const EDDY::CudaVolume& zcoord,
unsigned int xzp,
bool sync,
thrust::device_vector<float>& Wir) const EddyTry
{
int tpb = _threads_per_block_Wir;
EddyMatrixKernels::Wir<<<zcoord.Size(0),tpb>>>(zcoord.GetPtr(),zcoord.Size(0),zcoord.Size(1),
zcoord.Size(2),zcoord.Size(0),xzp,
thrust::raw_pointer_cast(Wir.data()));
if (sync) EddyCudaHelperFunctions::CudaSync("EddyMatrixKernels::Wir");
} EddyCatch
void StackResampler::make_Wir_t_y_vectors(const EDDY::CudaVolume& y,
const thrust::device_vector<float>& Wir,
unsigned int xzp,
bool sync,
thrust::device_vector<float>& Wirty) const EddyTry
{
int tpb = _threads_per_block_Wirty;
EddyMatrixKernels::Wirty<<<y.Size(0),tpb>>>(y.GetPtr(),thrust::raw_pointer_cast(Wir.data()),
y.Size(0),y.Size(1),y.Size(2),y.Size(0),xzp,
thrust::raw_pointer_cast(Wirty.data()));
if (sync) EddyCudaHelperFunctions::CudaSync("EddyMatrixKernels::Wirty");
} EddyCatch
void StackResampler::make_WtW_StS_matrices(const thrust::device_vector<float>& Wir,
unsigned int mn,
unsigned int nmat,
const thrust::device_vector<float>& StS,
bool sync,
thrust::device_vector<float>& WtW) const EddyTry
{
dim3 block = _threads_per_block_WtW_StS;
EddyMatrixKernels::KtK<<<nmat,block>>>(thrust::raw_pointer_cast(Wir.data()),mn,mn,nmat,
thrust::raw_pointer_cast(StS.data()),1.0,true,
thrust::raw_pointer_cast(WtW.data()));
if (sync) EddyCudaHelperFunctions::CudaSync("KtK_Kernels::KtK");
return;
} EddyCatch
void StackResampler::solve_for_c_hat(// Input
const thrust::device_vector<float>& WtW, // WtW+StS matrices for one xz-plane.
const thrust::device_vector<float>& Wty, // Wty vectors for one xz-plane
unsigned int n, // Size of KtK (nxn)
unsigned int nmat, // Number of matrices for one xz-plane
bool sync, // If true syncs after submitting kernel
// Output
thrust::device_vector<float>& chat) const EddyTry // Returns inv(WtW)*Wty for all matrices in xz-plane
{
// Allocate memory for Q and R matrices for QR decomposition
thrust::device_vector<float> Qt(nmat*n*n);
thrust::device_vector<float> R(nmat*n*n);
// Dynamically allocated shared memory, per block (matrix)
size_t sh_mem_sz = 2*n*sizeof(float);
int tpb = _threads_per_block_QR;
EddyMatrixKernels::QR<<<nmat,tpb,sh_mem_sz>>>(thrust::raw_pointer_cast(WtW.data()),n,n,nmat,
thrust::raw_pointer_cast(Qt.data()),
thrust::raw_pointer_cast(R.data()));
if (sync) EddyCudaHelperFunctions::CudaSync("QR_Kernels::QR");
tpb = _threads_per_block_Solve;
EddyMatrixKernels::Solve<<<nmat,tpb>>>(thrust::raw_pointer_cast(Qt.data()),
thrust::raw_pointer_cast(R.data()),
thrust::raw_pointer_cast(Wty.data()),n,n,nmat,
thrust::raw_pointer_cast(chat.data()));
if (sync) EddyCudaHelperFunctions::CudaSync("QR_Kernels::Solve");
return;
} EddyCatch
void StackResampler::make_y_hat_vectors(// Input
const thrust::device_vector<float>& W,
const thrust::device_vector<float>& chat,
unsigned int mn,
unsigned int nvec,
bool sync,
// Output
thrust::device_vector<float>& yhat) const EddyTry
{
int tpb = _threads_per_block_yhat;
EddyMatrixKernels::Ab<<<nvec,tpb>>>(thrust::raw_pointer_cast(W.data()),
thrust::raw_pointer_cast(chat.data()),
mn,mn,1,nvec,thrust::raw_pointer_cast(yhat.data()));
if (sync) EddyCudaHelperFunctions::CudaSync("EddyMatrixKernels::Ab");
return;
} EddyCatch
void StackResampler::transfer_y_hat_to_volume(// Input
const thrust::device_vector<float>& yhat,
unsigned int xzp,
bool sync,
// Output
EDDY::CudaVolume& ovol) const EddyTry
{
int tpb = _threads_per_block_transfer;
int nblocks = (ovol.Size(0)%tpb) ? ovol.Size(0) / tpb + 1 : ovol.Size(0) / tpb;
EddyKernels::transfer_y_hat_to_volume<<<nblocks,tpb>>>(thrust::raw_pointer_cast(yhat.data()),ovol.Size(0),
ovol.Size(1),ovol.Size(2),xzp,ovol.GetPtr());
if (sync) EddyCudaHelperFunctions::CudaSync("EddyMatrixKernels::transfer_y_hat_to_volume");
return;
} EddyCatch
void StackResampler::write_matrix(const thrust::device_vector<float>& mats,
unsigned int offs,
unsigned int m,
unsigned int n,
const std::string& fname) const EddyTry
{
thrust::device_vector<float>::const_iterator first = mats.begin() + offs;
thrust::device_vector<float>::const_iterator last = mats.begin() + offs + m*n;
thrust::host_vector<float> mat(first,last);
NEWMAT::Matrix newmat(m,n);
for (unsigned int i=0; i<m; i++) {
for (unsigned int j=0; j<n; j++) {
newmat(i+1,j+1) = mat[rfindx(i,j,m)];
}
}
MISCMATHS::write_ascii_matrix(fname+std::string(".txt"),newmat);
} EddyCatch
void StackResampler::write_matrices(const thrust::device_vector<float>& mats,
unsigned int nmat,
unsigned int m,
unsigned int n,
const std::string& basefname) const EddyTry
{
char fname[256];
for (unsigned int f=0; f<nmat; f++) {
sprintf(fname,"%s_%03d",basefname.c_str(),f);
write_matrix(mats,f*m*n,m,n,std::string(fname));
}
} EddyCatch
void StackResampler::write_debug_info_for_pred_resampling(unsigned int x,
unsigned int y,
const std::string& bfname,
const EDDY::CudaVolume& z,
const EDDY::CudaVolume& g,
const EDDY::CudaVolume& p,
const thrust::device_vector<float>& sz,
const thrust::device_vector<float>& W,
const thrust::device_vector<float>& Wir,
const thrust::device_vector<float>& w,
const thrust::device_vector<float>& wp,
const thrust::device_vector<float>& wW,
const thrust::device_vector<float>& Wirtg,
const thrust::device_vector<float>& wWtwp,
const thrust::device_vector<float>& WirtWir,
const thrust::device_vector<float>& wWtwW,
const thrust::device_vector<float>& sum_vec,
const thrust::device_vector<float>& sum_mat,
const thrust::device_vector<float>& c_hat,
const thrust::device_vector<float>& y_hat) const EddyTry
{
unsigned int xs = z.Size(0);
unsigned int ys = z.Size(1);
unsigned int zs = z.Size(2);
NEWIMAGE::volume<float> tmpvol = z.GetVolume();
NEWMAT::ColumnVector tmpvec(zs);
for (unsigned int k=0; k<zs; k++) tmpvec(k+1) = tmpvol(x,y,k);
std::string tmpfname = bfname + "_z";
MISCMATHS::write_ascii_matrix(tmpfname,tmpvec);
tmpvol = g.GetVolume();
for (unsigned int k=0; k<zs; k++) tmpvec(k+1) = tmpvol(x,y,k);
tmpfname = bfname + "_g";
MISCMATHS::write_ascii_matrix(tmpfname,tmpvec);
tmpvol = p.GetVolume();
for (unsigned int k=0; k<zs; k++) tmpvec(k+1) = tmpvol(x,y,k);
tmpfname = bfname + "_p";
MISCMATHS::write_ascii_matrix(tmpfname,tmpvec);
tmpfname = bfname + "_sz";
write_matrix(sz,(y*xs+x)*zs,zs,1,tmpfname);
tmpfname = bfname + "_W";
write_matrix(W,0,zs,zs,tmpfname);
tmpfname = bfname + "_Wir";
write_matrix(Wir,x*zs*zs,zs,zs,tmpfname);
tmpfname = bfname + "_wgt";
write_matrix(w,x*zs,zs,1,tmpfname);
tmpfname = bfname + "_wp";
write_matrix(wp,x*zs,zs,1,tmpfname);
tmpfname = bfname + "_wW";
write_matrix(wW,x*zs*zs,zs,zs,tmpfname);
tmpfname = bfname + "_Wirtg";
write_matrix(Wirtg,x*zs,zs,1,tmpfname);
tmpfname = bfname + "_wWtwp";
write_matrix(wWtwp,x*zs,zs,1,tmpfname);
tmpfname = bfname + "_WirtWir";
write_matrix(WirtWir,x*zs*zs,zs,zs,tmpfname);
tmpfname = bfname + "_wWtwW";
write_matrix(wWtwW,x*zs*zs,zs,zs,tmpfname);
tmpfname = bfname + "_sum_vec";
write_matrix(sum_vec,x*zs,zs,1,tmpfname);
tmpfname = bfname + "_sum_mat";
write_matrix(sum_mat,x*zs*zs,zs,zs,tmpfname);
tmpfname = bfname + "_c_hat";
write_matrix(c_hat,x*zs,zs,1,tmpfname);
tmpfname = bfname + "_y_hat";
write_matrix(y_hat,x*zs,zs,1,tmpfname);
} EddyCatch
} // End namespace EDDY
cuda/StackResampler.cu.prehip 0 → 100644
View file @ b9e18ed6
/*! \file StackResampler.cu
\brief Contains definition of CUDA implementation of a class for spline resampling of irregularly sampled columns in the z-direction
\author Jesper Andersson
\version 1.0b, March, 2016.
*/
//
// StackResampler.cu
//
// Jesper Andersson, FMRIB Image Analysis Group
//
// Copyright (C) 2016 University of Oxford
//
// Because of a bug in cuda_fp16.hpp, that gets included by cublas_v2.h, it has to
// be included before any include files that set up anything related to the std-lib.
// If not, there will be an ambiguity in cuda_fp16.hpp about wether to use the
// old-style C isinf or the new (since C++11) std::isinf.
#include "cublas_v2.h"
#include <cstdlib>
#include <string>
#include <vector>
#include <cmath>
#include <cuda.h>
#include <thrust/system_error.h>
#include <thrust/host_vector.h>
#include <thrust/device_vector.h>
#include <thrust/device_ptr.h>
#include <thrust/transform.h>
#include <thrust/fill.h>
#pragma push
#pragma diag_suppress = code_is_unreachable // Supress warnings from armawrap
#include "armawrap/newmat.h"
#include "newimage/newimage.h"
#include "EddyHelperClasses.h"
#include "CudaVolume.h"
#pragma pop
#include "EddyCudaHelperFunctions.h"
#include "StackResampler.h"
#include "EddyKernels.h"
#include "EddyMatrixKernels.h"
#include "EddyFunctors.h"
namespace EDDY {
const dim3 StackResampler::_threads_per_block_WtW_StS = dim3(16,16);
StackResampler::StackResampler(const EDDY::CudaVolume& stack,
const EDDY::CudaVolume& zcoord,
const EDDY::CudaVolume& pred,
const EDDY::CudaVolume& mask,
double lambda) EddyTry : _resvol(stack,false), _mask(stack,false)
{
// static unsigned int cnt = 0;
unsigned int matsz = stack.Size(2);
unsigned int nmat = stack.Size(0); // Number of matrices per xz-plane
// Allocate memory for storing matrices/vectors on the GPU
thrust::device_vector<float> StS_matrix(sqr(matsz));
thrust::device_vector<float> empty_StS_matrix(sqr(matsz));
thrust::device_vector<float> gpu_W_matrix(sqr(matsz));
thrust::device_vector<float> gpu_sorted_zcoord(zcoord.Size());
thrust::device_vector<float> gpu_Wir_matrices(nmat*sqr(matsz));
thrust::device_vector<float> gpu_weights(nmat*matsz);
thrust::device_vector<float> gpu_diagw_p_vectors(nmat*matsz);
thrust::device_vector<float> gpu_diagw_W_matrices(nmat*sqr(matsz));
thrust::device_vector<float> gpu_Wirty_vectors(nmat*matsz);
thrust::device_vector<float> gpu_dWtdp_vectors(nmat*matsz);
thrust::device_vector<float> gpu_WtW_matrices(nmat*sqr(matsz));
thrust::device_vector<float> gpu_dWtdW_matrices(nmat*sqr(matsz));
thrust::device_vector<float> gpu_sum_vectors(nmat*matsz);
thrust::device_vector<float> gpu_sum_matrices(nmat*sqr(matsz));
thrust::device_vector<float> gpu_c_hat_vectors(nmat*matsz);
thrust::device_vector<float> gpu_y_hat_vectors(nmat*matsz);
// EDDY::CudaVolume weights = pred;
// weights = 0.0;
// Get regularisation matrix once and for all
//cout << "Making reg matrix" << endl; cout.flush();
get_StS(matsz,lambda,StS_matrix);
get_StS(matsz,0.0,empty_StS_matrix);
// Get weight matrix for regular grid once and for all
//cout << "Making W" << endl; cout.flush();
get_regular_W(matsz,gpu_W_matrix);
// Populate mask
//cout << "Making mask" << endl; cout.flush();
make_mask(mask,zcoord,true,_mask);
// Get a volume of sorted z-ccords
//cout << "sorting zcoords" << endl; cout.flush();
sort_zcoords(zcoord,true,gpu_sorted_zcoord);
// Interpolate along z-columns
for (unsigned int j=0; j<stack.Size(1); j++) { // Loop over all xz-planes
// Make all Wir matrices for this xz-plane
//cout << "Making Wir" << endl; cout.flush();
make_Wir_matrices(zcoord,j,true,gpu_Wir_matrices);
// Make vectors of weights for predictions
//cout << "Making weights" << endl; cout.flush();
make_weight_vectors(gpu_sorted_zcoord,stack.Size(0),stack.Size(2),j,true,gpu_weights);
// insert_weights(gpu_weights,j,true,weights);
// Make diag{w}p vectors, where p are the predictions
//cout << "Making diag{w}p" << endl; cout.flush();
make_diagw_p_vectors(pred,gpu_weights,j,true,gpu_diagw_p_vectors);
// Make diag{w}W matrices, where W is the spline design matrix for regular sampling
//cout << "Making diag{w}W" << endl; cout.flush();
make_diagw_W_matrices(gpu_weights,gpu_W_matrix,matsz,nmat,true,gpu_diagw_W_matrices);
// Calculate all Wir'*y vectors for this plane
//cout << "Making Wirty" << endl; cout.flush();
make_Wir_t_y_vectors(stack,gpu_Wir_matrices,j,true,gpu_Wirty_vectors);
// Calculate all (diag{w}W)'*diag{w}p vectors for this plane
//cout << "Making diag{w}W'diag{w}p" << endl; cout.flush();
make_dwWt_dwp_vectors(gpu_diagw_W_matrices,gpu_diagw_p_vectors,matsz,nmat,true,gpu_dWtdp_vectors);
// Calculate all Wir'*Wir matrices
//cout << "Making Wir'Wir" << endl; cout.flush();
make_WtW_StS_matrices(gpu_Wir_matrices,matsz,nmat,StS_matrix,true,gpu_WtW_matrices);
// Calculate all (diag{w}W)'*(diag{w}W) matrices
//cout << "Making diag{w}W'diag{w}W" << endl; cout.flush();
make_WtW_StS_matrices(gpu_diagw_W_matrices,matsz,nmat,empty_StS_matrix,true,gpu_dWtdW_matrices);
// Add (pairwise) all Wir'*y and (diag{w}W)'*diag{w}p vectors
//cout << "Making sum vectors" << endl; cout.flush();
thrust::transform(gpu_Wirty_vectors.begin(),gpu_Wirty_vectors.end(),gpu_dWtdp_vectors.begin(),gpu_sum_vectors.begin(),thrust::plus<float>());
// Add (pairwise) all Wir'*Wir and (diag{w}W)'*(diag{w}W)
//cout << "Making sum matrices" << endl; cout.flush();
thrust::transform(gpu_WtW_matrices.begin(),gpu_WtW_matrices.end(),gpu_dWtdW_matrices.begin(),gpu_sum_matrices.begin(),thrust::plus<float>());
// Solve for spline coefficients
//cout << "Making c_hat" << endl; cout.flush();
solve_for_c_hat(gpu_sum_matrices,gpu_sum_vectors,matsz,nmat,true,gpu_c_hat_vectors);
// Multiply by W to solve for y_hat (interpolate in z-direction)
//cout << "Making y_hat" << endl; cout.flush();
make_y_hat_vectors(gpu_W_matrix,gpu_c_hat_vectors,matsz,nmat,true,gpu_y_hat_vectors);
// Transfer solution vectors to result volume
//cout << "Transferring y_hat" << endl; cout.flush();
transfer_y_hat_to_volume(gpu_y_hat_vectors,j,true,_resvol);
/*
if (cnt==19 && j==30) {
write_debug_info_for_pred_resampling(54,j,"debug_info.txt",zcoord,stack,pred,gpu_sorted_zcoord,
gpu_W_matrix,gpu_Wir_matrices,gpu_weights,gpu_diagw_p_vectors,
gpu_diagw_W_matrices,gpu_Wirty_vectors,gpu_dWtdp_vectors,
gpu_WtW_matrices,gpu_dWtdW_matrices,gpu_sum_vectors,
gpu_sum_matrices,gpu_c_hat_vectors,gpu_y_hat_vectors);
}
*/
}
// char fname[256]; sprintf(fname,"weights_%02d",cnt);
// weights.Write(std::string(fname));
// cnt++;
} EddyCatch
StackResampler::StackResampler(const EDDY::CudaVolume& stack,
const EDDY::CudaVolume& zcoord,
const EDDY::CudaVolume& mask,
NEWIMAGE::interpolation interp,
double lambda) EddyTry : _resvol(stack,false), _mask(stack,false)
{
// Populate mask
make_mask(mask,zcoord,true,_mask);
if (interp == NEWIMAGE::spline) {
unsigned int matsz = stack.Size(2);
unsigned int nmat = stack.Size(0); // Number of matrices per xz-plane
// Allocate memory for storing matrices/vectors on the GPU
thrust::device_vector<float> gpu_StS_matrix(sqr(matsz));
thrust::device_vector<float> gpu_W_matrix(sqr(matsz));
thrust::device_vector<float> gpu_Wir_matrices(nmat*sqr(matsz));
thrust::device_vector<float> gpu_WtW_StS_matrices(nmat*sqr(matsz));
thrust::device_vector<float> gpu_Wirty_vectors(nmat*matsz);
thrust::device_vector<float> gpu_c_hat_vectors(nmat*matsz);
thrust::device_vector<float> gpu_y_hat_vectors(nmat*matsz);
// Get regularisation matrix once and for all
get_StS(matsz,lambda,gpu_StS_matrix);
// Get spline matrix for regular grid once and for all
get_regular_W(matsz,gpu_W_matrix);
// Interpolate along z-columns
for (unsigned int j=0; j<stack.Size(1); j++) { // Loop over all xz-planes
// Make all Wir matrices for this xz-plane
make_Wir_matrices(zcoord,j,true,gpu_Wir_matrices);
// Pre-multiply Wir by transpose of self and add lambda*StS.
make_WtW_StS_matrices(gpu_Wir_matrices,matsz,nmat,gpu_StS_matrix,true,gpu_WtW_StS_matrices);
// Multiply Wir transposed with y (observed intensities).
make_Wir_t_y_vectors(stack,gpu_Wir_matrices,j,true,gpu_Wirty_vectors);
// Solve for spline coefficients
solve_for_c_hat(gpu_WtW_StS_matrices,gpu_Wirty_vectors,matsz,nmat,true,gpu_c_hat_vectors);
// Multiply by W to solve for y_hat (interpolate in z-direction)
make_y_hat_vectors(gpu_W_matrix,gpu_c_hat_vectors,matsz,nmat,true,gpu_y_hat_vectors);
// Transfer solution vectors to result volume
transfer_y_hat_to_volume(gpu_y_hat_vectors,j,true,_resvol);
}
}
else if (interp == NEWIMAGE::trilinear) {
thrust::device_vector<float> gpu_sorted_zcoord(zcoord.Size());
thrust::device_vector<float> gpu_sorted_intensities(zcoord.Size());
thrust::device_vector<float> gpu_interpolated_columns(zcoord.Size(),0.0);
// Sort the z-ccordinates and the intensities with z-coord as key
sort_zcoords_and_intensities(zcoord,stack,true,gpu_sorted_zcoord,gpu_sorted_intensities);
// Do linear interpolation along z-columns
linear_interpolate_columns(gpu_sorted_zcoord,gpu_sorted_intensities,zcoord.Size(0),zcoord.Size(1),zcoord.Size(2),true,gpu_interpolated_columns);
// Transfer interpolated vectors to results volume
transfer_interpolated_columns_to_volume(gpu_interpolated_columns,true,_resvol);
}
else throw EddyException("StackResampler::StackResampler: Invalid interpolation method");
} EddyCatch
void StackResampler::get_StS(unsigned int sz, double lambda, thrust::device_vector<float>& StS) const EddyTry
{
float six = lambda * 6.0; float minusfour = - lambda * 4.0; float one = lambda;
thrust::host_vector<float> hStS(StS.size(),0.0);
hStS[rfindx(0,0,sz)] = six; hStS[rfindx(0,1,sz)] = minusfour; hStS[rfindx(0,2,sz)] = one; hStS[rfindx(0,(sz-2),sz)] = one; hStS[rfindx(0,(sz-1),sz)] = minusfour;
hStS[rfindx(1,0,sz)] = minusfour; hStS[rfindx(1,1,sz)] = six; hStS[rfindx(1,2,sz)] = minusfour; hStS[rfindx(1,3,sz)] = one; hStS[rfindx(1,(sz-1),sz)] = one;
for ( unsigned int i=2; i<(sz-2); i++) {
hStS[rfindx(i,i-2,sz)] = one;
hStS[rfindx(i,i-1,sz)] = minusfour;
hStS[rfindx(i,i,sz)] = six;
hStS[rfindx(i,i+1,sz)] = minusfour;
hStS[rfindx(i,i+2,sz)] = one;
}
hStS[rfindx((sz-2),0,sz)] = one; hStS[rfindx((sz-2),(sz-4),sz)] = one; hStS[rfindx((sz-2),(sz-3),sz)] = minusfour; hStS[rfindx((sz-2),(sz-2),sz)] = six; hStS[rfindx((sz-2),(sz-1),sz)] = minusfour;
hStS[rfindx((sz-1),0,sz)] = minusfour; hStS[rfindx((sz-1),1,sz)] = one; hStS[rfindx((sz-1),(sz-3),sz)] = one; hStS[rfindx((sz-1),(sz-2),sz)] = minusfour; hStS[rfindx((sz-1),(sz-1),sz)] = six;
StS = hStS;
return;
} EddyCatch
void StackResampler::get_regular_W(unsigned int sz, thrust::device_vector<float>& W) const EddyTry
{
thrust::host_vector<float> hW(W.size(),0.0);
hW[rfindx(0,0,sz)] = 5.0/6.0; hW[rfindx(0,1,sz)] = 1.0/6.0;
for ( unsigned int i=1; i<(sz-1); i++) {
hW[rfindx(i,i-1,sz)] = 1.0/6.0;
hW[rfindx(i,i,sz)] = 4.0/6.0;
hW[rfindx(i,i+1,sz)] = 1.0/6.0;
}
hW[rfindx((sz-1),(sz-2),sz)] = 1.0/6.0; hW[rfindx((sz-1),(sz-1),sz)] = 5.0/6.0;
W = hW;
return;
} EddyCatch
void StackResampler::make_mask(const EDDY::CudaVolume& inmask,
const EDDY::CudaVolume& zcoord,
bool sync,
EDDY::CudaVolume& omask) EddyTry
{
omask = 0.0;
int nblocks = inmask.Size(1);
int tpb = inmask.Size(0);
EddyKernels::make_mask_from_stack<<<nblocks,tpb>>>(inmask.GetPtr(),zcoord.GetPtr(),inmask.Size(0),
inmask.Size(1),inmask.Size(2),omask.GetPtr());
if (sync) EddyCudaHelperFunctions::CudaSync("EddyKernels::make_mask_from_stack");
} EddyCatch
void StackResampler::sort_zcoords(const EDDY::CudaVolume& zcoord,
bool sync,
thrust::device_vector<float>& szcoord) const EddyTry
{
// First transfer all columns and check if they are sorted
thrust::device_vector<unsigned int> ns_flags(zcoord.Size(0)*zcoord.Size(1),0); // Non-zero if a column is "not sorted"
EddyKernels::TransferAndCheckSorting<<<zcoord.Size(1),zcoord.Size(0)>>>(zcoord.GetPtr(),zcoord.Size(0),
zcoord.Size(1),zcoord.Size(2),
thrust::raw_pointer_cast(szcoord.data()),
thrust::raw_pointer_cast(ns_flags.data()));
EddyCudaHelperFunctions::CudaSync("EddyKernels::TransferAndCheckSorting");
unsigned int nnsort = thrust::reduce(ns_flags.begin(),ns_flags.end(),0,EDDY::Sum<unsigned int,unsigned int>());
if (nnsort) { // If there are columns that are not sorted
thrust::host_vector<unsigned int> host_ns_flags = ns_flags;
thrust::host_vector<unsigned int> host_nsort_indx(nnsort,0);
for (unsigned int i=0, n=0; i<zcoord.Size(0)*zcoord.Size(1); i++) {
if (host_ns_flags[i]) { host_nsort_indx[n] = i; n++; }
}
thrust::device_vector<unsigned int> device_nsort_indx = host_nsort_indx;
int nb = (nnsort / 32) + 1;
EddyKernels::SortVectors<<<nb,32>>>(thrust::raw_pointer_cast(device_nsort_indx.data()),nnsort,
zcoord.Size(2),thrust::raw_pointer_cast(szcoord.data()),NULL);
if (sync) EddyCudaHelperFunctions::CudaSync("EddyKernels::SortVectors");
/*
// Check that vectors really are sorted. Only for use during testing.
thrust::host_vector<float> hvec = szcoord;
for (unsigned int vi=0; vi<zcoord.Size(1)*zcoord.Size(0); vi++) {
unsigned int offs = vi*zcoord.Size(2);
for (unsigned int i=1; i<zcoord.Size(2); i++) {
if (hvec[offs+i] < hvec[offs+i-1]) {
cout << "Vector " << vi << " was not sucessfully sorted" << endl;
cout.flush();
for (unsigned int j=0; j<zcoord.Size(2); j++) {
cout << "hvec[" << j << "] = " << hvec[offs+j] << endl;
}
exit(0);
}
}
}
*/
}
} EddyCatch
void StackResampler::sort_zcoords_and_intensities(const EDDY::CudaVolume& zcoord,
const EDDY::CudaVolume& data,
bool sync,
thrust::device_vector<float>& szcoord,
thrust::device_vector<float>& sdata) const EddyTry
{
// First transfer all zcoord columns and check if they are sorted
thrust::device_vector<unsigned int> ns_flags(zcoord.Size(0)*zcoord.Size(1),0); // Non-zero if a column is "not sorted"
EddyKernels::TransferAndCheckSorting<<<zcoord.Size(1),zcoord.Size(0)>>>(zcoord.GetPtr(),zcoord.Size(0),
zcoord.Size(1),zcoord.Size(2),
thrust::raw_pointer_cast(szcoord.data()),
thrust::raw_pointer_cast(ns_flags.data()));
EddyCudaHelperFunctions::CudaSync("EddyKernels::TransferAndCheckSorting");
unsigned int nnsort = thrust::reduce(ns_flags.begin(),ns_flags.end(),0,EDDY::Sum<unsigned int,unsigned int>());
// Transfer all intensity values as well
EddyKernels::TransferVolumeToVectors<<<data.Size(1),data.Size(0)>>>(data.GetPtr(),data.Size(0),
data.Size(1),data.Size(2),
thrust::raw_pointer_cast(sdata.data()));
if (nnsort) { // If there are columns that are not sorted
thrust::host_vector<unsigned int> host_ns_flags = ns_flags;
thrust::host_vector<unsigned int> host_nsort_indx(nnsort,0);
for (unsigned int i=0, n=0; i<zcoord.Size(0)*zcoord.Size(1); i++) {
if (host_ns_flags[i]) { host_nsort_indx[n] = i; n++; }
}
thrust::device_vector<unsigned int> device_nsort_indx = host_nsort_indx;
int nb = (nnsort / 32) + 1;
EddyKernels::SortVectors<<<nb,32>>>(thrust::raw_pointer_cast(device_nsort_indx.data()),nnsort,
zcoord.Size(2),thrust::raw_pointer_cast(szcoord.data()),
thrust::raw_pointer_cast(sdata.data()));
if (sync) EddyCudaHelperFunctions::CudaSync("EddyKernels::SortVectors");
}
} EddyCatch
void StackResampler::linear_interpolate_columns(const thrust::device_vector<float>& zcoord,
const thrust::device_vector<float>& val,
unsigned int xsz,
unsigned int ysz,
unsigned int zsz,
bool sync,
thrust::device_vector<float>& ival) const EddyTry
{
EddyKernels::LinearInterpolate<<<ysz,xsz>>>(thrust::raw_pointer_cast(zcoord.data()),
thrust::raw_pointer_cast(val.data()),zsz,
thrust::raw_pointer_cast(ival.data()));
if (sync) EddyCudaHelperFunctions::CudaSync("EddyKernels::LinearInterpolate");
} EddyCatch
void StackResampler::transfer_interpolated_columns_to_volume(const thrust::device_vector<float>& zcols,
bool sync,
EDDY::CudaVolume& vol) EddyTry
{
EddyKernels::TransferColumnsToVolume<<<vol.Size(1),vol.Size(0)>>>(thrust::raw_pointer_cast(zcols.data()),
vol.Size(0),vol.Size(1),
vol.Size(2),vol.GetPtr());
if (sync) EddyCudaHelperFunctions::CudaSync("EddyKernels::TransferColumnsToVolume");
} EddyCatch
void StackResampler::make_weight_vectors(const thrust::device_vector<float>& zcoord,
unsigned int xsz,
unsigned int zsz,
unsigned int xzp,
bool sync,
thrust::device_vector<float>& weights) const EddyTry
{
EddyKernels::MakeWeights<<<xsz,zsz>>>(thrust::raw_pointer_cast(zcoord.data()),xsz,
zsz,xzp,thrust::raw_pointer_cast(weights.data()));
if (sync) EddyCudaHelperFunctions::CudaSync("EddyKernels::MakeWeights");
} EddyCatch
void StackResampler::insert_weights(const thrust::device_vector<float>& wvec,
unsigned int j,
bool sync,
EDDY::CudaVolume& wvol) const EddyTry
{
EddyKernels::InsertWeights<<<wvol.Size(0),wvol.Size(2)>>>(thrust::raw_pointer_cast(wvec.data()),j,wvol.Size(0),
wvol.Size(1),wvol.Size(2),wvol.GetPtr());
if (sync) EddyCudaHelperFunctions::CudaSync("EddyKernels::InsertWeights");
} EddyCatch
void StackResampler::make_diagw_p_vectors(const EDDY::CudaVolume& pred,
const thrust::device_vector<float>& wgts,
unsigned int xzp,
bool sync,
thrust::device_vector<float>& wp) const EddyTry
{
EddyKernels::MakeDiagwpVecs<<<pred.Size(0),pred.Size(2)>>>(pred.GetPtr(),thrust::raw_pointer_cast(wgts.data()),
pred.Size(0),pred.Size(1),pred.Size(2),xzp,
thrust::raw_pointer_cast(wp.data()));
if (sync) EddyCudaHelperFunctions::CudaSync("EddyKernels::MakeDiagwpVecs");
} EddyCatch
void StackResampler::make_diagw_W_matrices(const thrust::device_vector<float>& wgts,
const thrust::device_vector<float>& W,
unsigned int matsz,
unsigned int nmat,
bool sync,
thrust::device_vector<float>& diagwW) const EddyTry
{
EddyMatrixKernels::DiagwA<<<nmat,matsz>>>(thrust::raw_pointer_cast(wgts.data()),
thrust::raw_pointer_cast(W.data()),
matsz,matsz,nmat,
thrust::raw_pointer_cast(diagwW.data()));
if (sync) EddyCudaHelperFunctions::CudaSync("EddyMatrixKernels::DiagwA");
} EddyCatch
void StackResampler::make_dwWt_dwp_vectors(const thrust::device_vector<float>& dW,
const thrust::device_vector<float>& dp,
unsigned int matsz,
unsigned int nmat,
bool sync,
thrust::device_vector<float>& dWtdp) const EddyTry
{
EddyMatrixKernels::Atb<<<nmat,matsz>>>(thrust::raw_pointer_cast(dW.data()),
thrust::raw_pointer_cast(dp.data()),
matsz,matsz,nmat,nmat,
thrust::raw_pointer_cast(dWtdp.data()));
if (sync) EddyCudaHelperFunctions::CudaSync("EddyMatrixKernels::Atb");
return;
} EddyCatch
void StackResampler::make_Wir_matrices(const EDDY::CudaVolume& zcoord,
unsigned int xzp,
bool sync,
thrust::device_vector<float>& Wir) const EddyTry
{
int tpb = _threads_per_block_Wir;
EddyMatrixKernels::Wir<<<zcoord.Size(0),tpb>>>(zcoord.GetPtr(),zcoord.Size(0),zcoord.Size(1),
zcoord.Size(2),zcoord.Size(0),xzp,
thrust::raw_pointer_cast(Wir.data()));
if (sync) EddyCudaHelperFunctions::CudaSync("EddyMatrixKernels::Wir");
} EddyCatch
void StackResampler::make_Wir_t_y_vectors(const EDDY::CudaVolume& y,
const thrust::device_vector<float>& Wir,
unsigned int xzp,
bool sync,
thrust::device_vector<float>& Wirty) const EddyTry
{
int tpb = _threads_per_block_Wirty;
EddyMatrixKernels::Wirty<<<y.Size(0),tpb>>>(y.GetPtr(),thrust::raw_pointer_cast(Wir.data()),
y.Size(0),y.Size(1),y.Size(2),y.Size(0),xzp,
thrust::raw_pointer_cast(Wirty.data()));
if (sync) EddyCudaHelperFunctions::CudaSync("EddyMatrixKernels::Wirty");
} EddyCatch
void StackResampler::make_WtW_StS_matrices(const thrust::device_vector<float>& Wir,
unsigned int mn,
unsigned int nmat,
const thrust::device_vector<float>& StS,
bool sync,
thrust::device_vector<float>& WtW) const EddyTry
{
dim3 block = _threads_per_block_WtW_StS;
EddyMatrixKernels::KtK<<<nmat,block>>>(thrust::raw_pointer_cast(Wir.data()),mn,mn,nmat,
thrust::raw_pointer_cast(StS.data()),1.0,true,
thrust::raw_pointer_cast(WtW.data()));
if (sync) EddyCudaHelperFunctions::CudaSync("KtK_Kernels::KtK");
return;
} EddyCatch
void StackResampler::solve_for_c_hat(// Input
const thrust::device_vector<float>& WtW, // WtW+StS matrices for one xz-plane.
const thrust::device_vector<float>& Wty, // Wty vectors for one xz-plane
unsigned int n, // Size of KtK (nxn)
unsigned int nmat, // Number of matrices for one xz-plane
bool sync, // If true syncs after submitting kernel
// Output
thrust::device_vector<float>& chat) const EddyTry // Returns inv(WtW)*Wty for all matrices in xz-plane
{
// Allocate memory for Q and R matrices for QR decomposition
thrust::device_vector<float> Qt(nmat*n*n);
thrust::device_vector<float> R(nmat*n*n);
// Dynamically allocated shared memory, per block (matrix)
size_t sh_mem_sz = 2*n*sizeof(float);
int tpb = _threads_per_block_QR;
EddyMatrixKernels::QR<<<nmat,tpb,sh_mem_sz>>>(thrust::raw_pointer_cast(WtW.data()),n,n,nmat,
thrust::raw_pointer_cast(Qt.data()),
thrust::raw_pointer_cast(R.data()));
if (sync) EddyCudaHelperFunctions::CudaSync("QR_Kernels::QR");
tpb = _threads_per_block_Solve;
EddyMatrixKernels::Solve<<<nmat,tpb>>>(thrust::raw_pointer_cast(Qt.data()),
thrust::raw_pointer_cast(R.data()),
thrust::raw_pointer_cast(Wty.data()),n,n,nmat,
thrust::raw_pointer_cast(chat.data()));
if (sync) EddyCudaHelperFunctions::CudaSync("QR_Kernels::Solve");
return;
} EddyCatch
void StackResampler::make_y_hat_vectors(// Input
const thrust::device_vector<float>& W,
const thrust::device_vector<float>& chat,
unsigned int mn,
unsigned int nvec,
bool sync,
// Output
thrust::device_vector<float>& yhat) const EddyTry
{
int tpb = _threads_per_block_yhat;
EddyMatrixKernels::Ab<<<nvec,tpb>>>(thrust::raw_pointer_cast(W.data()),
thrust::raw_pointer_cast(chat.data()),
mn,mn,1,nvec,thrust::raw_pointer_cast(yhat.data()));
if (sync) EddyCudaHelperFunctions::CudaSync("EddyMatrixKernels::Ab");
return;
} EddyCatch
void StackResampler::transfer_y_hat_to_volume(// Input
const thrust::device_vector<float>& yhat,
unsigned int xzp,
bool sync,
// Output
EDDY::CudaVolume& ovol) const EddyTry
{
int tpb = _threads_per_block_transfer;
int nblocks = (ovol.Size(0)%tpb) ? ovol.Size(0) / tpb + 1 : ovol.Size(0) / tpb;
EddyKernels::transfer_y_hat_to_volume<<<nblocks,tpb>>>(thrust::raw_pointer_cast(yhat.data()),ovol.Size(0),
ovol.Size(1),ovol.Size(2),xzp,ovol.GetPtr());
if (sync) EddyCudaHelperFunctions::CudaSync("EddyMatrixKernels::transfer_y_hat_to_volume");
return;
} EddyCatch
void StackResampler::write_matrix(const thrust::device_vector<float>& mats,
unsigned int offs,
unsigned int m,
unsigned int n,
const std::string& fname) const EddyTry
{
thrust::device_vector<float>::const_iterator first = mats.begin() + offs;
thrust::device_vector<float>::const_iterator last = mats.begin() + offs + m*n;
thrust::host_vector<float> mat(first,last);
NEWMAT::Matrix newmat(m,n);
for (unsigned int i=0; i<m; i++) {
for (unsigned int j=0; j<n; j++) {
newmat(i+1,j+1) = mat[rfindx(i,j,m)];
}
}
MISCMATHS::write_ascii_matrix(fname+std::string(".txt"),newmat);
} EddyCatch
void StackResampler::write_matrices(const thrust::device_vector<float>& mats,
unsigned int nmat,
unsigned int m,
unsigned int n,
const std::string& basefname) const EddyTry
{
char fname[256];
for (unsigned int f=0; f<nmat; f++) {
sprintf(fname,"%s_%03d",basefname.c_str(),f);
write_matrix(mats,f*m*n,m,n,std::string(fname));
}
} EddyCatch
void StackResampler::write_debug_info_for_pred_resampling(unsigned int x,
unsigned int y,
const std::string& bfname,
const EDDY::CudaVolume& z,
const EDDY::CudaVolume& g,
const EDDY::CudaVolume& p,
const thrust::device_vector<float>& sz,
const thrust::device_vector<float>& W,
const thrust::device_vector<float>& Wir,
const thrust::device_vector<float>& w,
const thrust::device_vector<float>& wp,
const thrust::device_vector<float>& wW,
const thrust::device_vector<float>& Wirtg,
const thrust::device_vector<float>& wWtwp,
const thrust::device_vector<float>& WirtWir,
const thrust::device_vector<float>& wWtwW,
const thrust::device_vector<float>& sum_vec,
const thrust::device_vector<float>& sum_mat,
const thrust::device_vector<float>& c_hat,
const thrust::device_vector<float>& y_hat) const EddyTry
{
unsigned int xs = z.Size(0);
unsigned int ys = z.Size(1);
unsigned int zs = z.Size(2);
NEWIMAGE::volume<float> tmpvol = z.GetVolume();
NEWMAT::ColumnVector tmpvec(zs);
for (unsigned int k=0; k<zs; k++) tmpvec(k+1) = tmpvol(x,y,k);
std::string tmpfname = bfname + "_z";
MISCMATHS::write_ascii_matrix(tmpfname,tmpvec);
tmpvol = g.GetVolume();
for (unsigned int k=0; k<zs; k++) tmpvec(k+1) = tmpvol(x,y,k);
tmpfname = bfname + "_g";
MISCMATHS::write_ascii_matrix(tmpfname,tmpvec);
tmpvol = p.GetVolume();
for (unsigned int k=0; k<zs; k++) tmpvec(k+1) = tmpvol(x,y,k);
tmpfname = bfname + "_p";
MISCMATHS::write_ascii_matrix(tmpfname,tmpvec);
tmpfname = bfname + "_sz";
write_matrix(sz,(y*xs+x)*zs,zs,1,tmpfname);
tmpfname = bfname + "_W";
write_matrix(W,0,zs,zs,tmpfname);
tmpfname = bfname + "_Wir";
write_matrix(Wir,x*zs*zs,zs,zs,tmpfname);
tmpfname = bfname + "_wgt";
write_matrix(w,x*zs,zs,1,tmpfname);
tmpfname = bfname + "_wp";
write_matrix(wp,x*zs,zs,1,tmpfname);
tmpfname = bfname + "_wW";
write_matrix(wW,x*zs*zs,zs,zs,tmpfname);
tmpfname = bfname + "_Wirtg";
write_matrix(Wirtg,x*zs,zs,1,tmpfname);
tmpfname = bfname + "_wWtwp";
write_matrix(wWtwp,x*zs,zs,1,tmpfname);
tmpfname = bfname + "_WirtWir";
write_matrix(WirtWir,x*zs*zs,zs,zs,tmpfname);
tmpfname = bfname + "_wWtwW";
write_matrix(wWtwW,x*zs*zs,zs,zs,tmpfname);
tmpfname = bfname + "_sum_vec";
write_matrix(sum_vec,x*zs,zs,1,tmpfname);
tmpfname = bfname + "_sum_mat";
write_matrix(sum_mat,x*zs*zs,zs,zs,tmpfname);
tmpfname = bfname + "_c_hat";
write_matrix(c_hat,x*zs,zs,1,tmpfname);
tmpfname = bfname + "_y_hat";
write_matrix(y_hat,x*zs,zs,1,tmpfname);
} EddyCatch
} // End namespace EDDY
cuda/StackResampler.h 0 → 100644
View file @ b9e18ed6
/*! \file StackResampler.h
\brief Contains declaration of CUDA implementation of a class for spline resampling of irregularly sampled columns in the z-direction
\author Jesper Andersson
\version 1.0b, March, 2016.
*/
//
// StackResampler.h
//
// Jesper Andersson, FMRIB Image Analysis Group
//
// Copyright (C) 2016 University of Oxford
//
#include <cstdlib>
#include <string>
#include <vector>
#include <cmath>
#include <hip/hip_runtime.h>
#include <thrust/system_error.h>
#include <thrust/host_vector.h>
#include <thrust/device_vector.h>
#include <thrust/device_ptr.h>
#include <thrust/transform.h>
#include <thrust/fill.h>
#pragma push
#pragma diag_suppress = code_is_unreachable // Supress warnings from armawrap
#include "armawrap/newmat.h"
#include "newimage/newimage.h"
#include "EddyHelperClasses.h"
#include "CudaVolume.h"
#pragma pop
namespace EDDY {
class StackResampler
{
public:
/// Constructor. Performs the actual work so that when the object is created there is already a resampled image ready.
/// This version of the constructor uses a predicted volume and the laplacian for regularisation
StackResampler(const EDDY::CudaVolume& stack,
const EDDY::CudaVolume& zcoord,
const EDDY::CudaVolume& pred,
const EDDY::CudaVolume& mask,
double lambda=0.005);
/// This version of the constructor uses either splines and Laplacian regularisation or linear interpolation.
StackResampler(const EDDY::CudaVolume& stack,
const EDDY::CudaVolume& zcoord,
const EDDY::CudaVolume& mask,
NEWIMAGE::interpolation interp=NEWIMAGE::spline,
double lambda=0.005);
~StackResampler() {}
/// Returns the resampled volue
const EDDY::CudaVolume& GetResampledIma() const EddyTry { return(_resvol); } EddyCatch
NEWIMAGE::volume<float> GetResampledImaAsNEWIMAGE() const EddyTry { return(_resvol.GetVolume()); } EddyCatch
/// Returns binary mask to indicate valid voxels
const EDDY::CudaVolume& GetMask() const EddyTry { return(_mask); } EddyCatch
private:
static const int _threads_per_block_QR = 128;
static const int _threads_per_block_Solve = 128;
static const int _threads_per_block_Wirty = 128;
static const int _threads_per_block_Wir = 128;
static const int _threads_per_block_yhat = 128;
static const int _threads_per_block_transfer = 128;
static const dim3 _threads_per_block_WtW_StS;
EDDY::CudaVolume _resvol;
EDDY::CudaVolume _mask;
/// Get Laplacian for regularisation. Runs on CPU.
void get_StS(unsigned int sz, double lambda, thrust::device_vector<float>& StS) const;
/// Get "design matrix" for splines on regular grid. Runs on CPU.
void get_regular_W(unsigned int sz, thrust::device_vector<float>& W) const;
/// Make binary mask with one for valid voxels
void make_mask(const EDDY::CudaVolume& inmask,
const EDDY::CudaVolume& zcoord,
bool zync,
EDDY::CudaVolume& omask);
/// Take the z-columns in zccord, sort them and put them in z-x-y order in szcoord
void sort_zcoords(const EDDY::CudaVolume& zcoord,
bool sync,
thrust::device_vector<float>& szcoord) const;
/// Make z-first vectors of weights for the predictions
void make_weight_vectors(const thrust::device_vector<float>& zcoord,
unsigned int xsz,
unsigned int zsz,
unsigned int xzp,
bool sync,
thrust::device_vector<float>& weights) const;
/// Insert one xz-plane of weights into volume
void insert_weights(const thrust::device_vector<float>& wvec,
unsigned int j,
bool sync,
EDDY::CudaVolume& wvol) const;
/// Elementwise multiplication of weights with predictions.
void make_diagw_p_vectors(const EDDY::CudaVolume& pred,
const thrust::device_vector<float>& wgts,
unsigned int xzp,
bool sync,
thrust::device_vector<float>& wp) const;
/// Premultiply W matrix by diag{w}
void make_diagw_W_matrices(const thrust::device_vector<float>& w,
const thrust::device_vector<float>& W,
unsigned int matsz,
unsigned int nmat,
bool sync,
thrust::device_vector<float>& diagwW) const;
/// Multiply a set of (diag{w}W)' matrices with a bunch of diag{w}p vectors
void make_dwWt_dwp_vectors(const thrust::device_vector<float>& dW,
const thrust::device_vector<float>& dp,
unsigned int matsz,
unsigned int nmat,
bool sync,
thrust::device_vector<float>& dWtdp) const;
/// Make a set of "design matrices" for irregularly sampled splines. GPU.
void make_Wir_matrices(const EDDY::CudaVolume& zcoord,
unsigned int xzp,
bool sync,
thrust::device_vector<float>& Wir) const;
/// Multiply a set of Wir matrices with a set of intensity vectors. GPU.
void make_Wir_t_y_vectors(const EDDY::CudaVolume& y,
const thrust::device_vector<float>& Wir,
unsigned int xzp,
bool sync,
thrust::device_vector<float>& Wirty) const;
/// Multiply a set of Wir matrices by transpose of self and add regularisation matrix. GPU.
void make_WtW_StS_matrices(const thrust::device_vector<float>& Wir,
unsigned int mn,
unsigned int nmat,
const thrust::device_vector<float>& StS,
bool sync,
thrust::device_vector<float>& WtW) const;
/// Solve for spline coefficients
void solve_for_c_hat(// Input
const thrust::device_vector<float>& WtW, // WtW+StS matrices for one xz-plane.
const thrust::device_vector<float>& Wty, // Wty vectors for one xz-plane
unsigned int n, // Size of KtK (nxn)
unsigned int nmat, // Number of matrices for one xz-plane
bool sync, // If true syncs after submitting kernel
// Output
thrust::device_vector<float>& chat) const; // Returns inv(WtW)*Wty for all matrices in xz-plane
/// Multiply spline coefficients with regularly sampled spline matrix
void make_y_hat_vectors(// Input
const thrust::device_vector<float>& W,
const thrust::device_vector<float>& chat,
unsigned int mn,
unsigned int nvec,
bool sync,
// Output
thrust::device_vector<float>& yhat) const;
/// Transfer y_hat (interpolated data) to volume
void transfer_y_hat_to_volume(// Input
const thrust::device_vector<float>& yhat,
unsigned int xzp,
bool sync,
// Output
EDDY::CudaVolume& ovol) const;
/// Sorts with zccord as key and reorders data in the same way
void sort_zcoords_and_intensities(const EDDY::CudaVolume& zcoord,
const EDDY::CudaVolume& data,
bool sync,
thrust::device_vector<float>& szcoord,
thrust::device_vector<float>& sdata) const;
/// Takes sorted vectors of z-coords and values and linearly interpolates onto 0--(zsz-1)
void linear_interpolate_columns(const thrust::device_vector<float>& zcoord,
const thrust::device_vector<float>& val,
unsigned int xsz,
unsigned int ysz,
unsigned int zsz,
bool sync,
thrust::device_vector<float>& ival) const;
void transfer_interpolated_columns_to_volume(const thrust::device_vector<float>& zcols,
bool sync,
EDDY::CudaVolume& vol);
/// i,j->linear index for row-first square matrix. Matrix addressed M(0:mn-1,0:mn-1)
unsigned int rfindx(unsigned int i, unsigned int j, unsigned int mn) const { return(i+j*mn); }
/// i,j->linear index for column-first square matrix. Matrix addressed M(0:mn-1,0:mn-1)
unsigned int cfindx(unsigned int i, unsigned int j, unsigned int mn) const { return(i*mn+j); }
/// Local sqr, just to be safe
template <typename T> T sqr(T a) const { return(a*a); }
/// Writes a single matrix residing on the GPU as a Newmat text file
void write_matrix(const thrust::device_vector<float>& mats,
unsigned int offs,
unsigned int m,
unsigned int n,
const std::string& fname) const;
/// Writes a block of matrices residing on the GPU as a set of Newmat text files
void write_matrices(const thrust::device_vector<float>& mats,
unsigned int nmat,
unsigned int m,
unsigned int n,
const std::string& basefname) const;
/// Writes all neccessary info for one column for debugging
void write_debug_info_for_pred_resampling(unsigned int x,
unsigned int y,
const std::string& bfname,
const EDDY::CudaVolume& z,
const EDDY::CudaVolume& g,
const EDDY::CudaVolume& p,
const thrust::device_vector<float>& sz,
const thrust::device_vector<float>& W,
const thrust::device_vector<float>& Wir,
const thrust::device_vector<float>& w,
const thrust::device_vector<float>& wp,
const thrust::device_vector<float>& wW,
const thrust::device_vector<float>& Wirtg,
const thrust::device_vector<float>& wWtwp,
const thrust::device_vector<float>& WirtWir,
const thrust::device_vector<float>& wWtwW,
const thrust::device_vector<float>& sum_vec,
const thrust::device_vector<float>& sum_mat,
const thrust::device_vector<float>& c_hat,
const thrust::device_vector<float>& y_hat) const;
};
} // End namespace EDDY
cuda/StackResampler.h.prehip 0 → 100644
View file @ b9e18ed6
/*! \file StackResampler.h
\brief Contains declaration of CUDA implementation of a class for spline resampling of irregularly sampled columns in the z-direction
\author Jesper Andersson
\version 1.0b, March, 2016.
*/
//
// StackResampler.h
//
// Jesper Andersson, FMRIB Image Analysis Group
//
// Copyright (C) 2016 University of Oxford
//
#include <cstdlib>
#include <string>
#include <vector>
#include <cmath>
#include <cuda.h>
#include <thrust/system_error.h>
#include <thrust/host_vector.h>
#include <thrust/device_vector.h>
#include <thrust/device_ptr.h>
#include <thrust/transform.h>
#include <thrust/fill.h>
#pragma push
#pragma diag_suppress = code_is_unreachable // Supress warnings from armawrap
#include "armawrap/newmat.h"
#include "newimage/newimage.h"
#include "EddyHelperClasses.h"
#include "CudaVolume.h"
#pragma pop
namespace EDDY {
class StackResampler
{
public:
/// Constructor. Performs the actual work so that when the object is created there is already a resampled image ready.
/// This version of the constructor uses a predicted volume and the laplacian for regularisation
StackResampler(const EDDY::CudaVolume& stack,
const EDDY::CudaVolume& zcoord,
const EDDY::CudaVolume& pred,
const EDDY::CudaVolume& mask,
double lambda=0.005);
/// This version of the constructor uses either splines and Laplacian regularisation or linear interpolation.
StackResampler(const EDDY::CudaVolume& stack,
const EDDY::CudaVolume& zcoord,
const EDDY::CudaVolume& mask,
NEWIMAGE::interpolation interp=NEWIMAGE::spline,
double lambda=0.005);
~StackResampler() {}
/// Returns the resampled volue
const EDDY::CudaVolume& GetResampledIma() const EddyTry { return(_resvol); } EddyCatch
NEWIMAGE::volume<float> GetResampledImaAsNEWIMAGE() const EddyTry { return(_resvol.GetVolume()); } EddyCatch
/// Returns binary mask to indicate valid voxels
const EDDY::CudaVolume& GetMask() const EddyTry { return(_mask); } EddyCatch
private:
static const int _threads_per_block_QR = 128;
static const int _threads_per_block_Solve = 128;
static const int _threads_per_block_Wirty = 128;
static const int _threads_per_block_Wir = 128;
static const int _threads_per_block_yhat = 128;
static const int _threads_per_block_transfer = 128;
static const dim3 _threads_per_block_WtW_StS;
EDDY::CudaVolume _resvol;
EDDY::CudaVolume _mask;
/// Get Laplacian for regularisation. Runs on CPU.
void get_StS(unsigned int sz, double lambda, thrust::device_vector<float>& StS) const;
/// Get "design matrix" for splines on regular grid. Runs on CPU.
void get_regular_W(unsigned int sz, thrust::device_vector<float>& W) const;
/// Make binary mask with one for valid voxels
void make_mask(const EDDY::CudaVolume& inmask,
const EDDY::CudaVolume& zcoord,
bool zync,
EDDY::CudaVolume& omask);
/// Take the z-columns in zccord, sort them and put them in z-x-y order in szcoord
void sort_zcoords(const EDDY::CudaVolume& zcoord,
bool sync,
thrust::device_vector<float>& szcoord) const;
/// Make z-first vectors of weights for the predictions
void make_weight_vectors(const thrust::device_vector<float>& zcoord,
unsigned int xsz,
unsigned int zsz,
unsigned int xzp,
bool sync,
thrust::device_vector<float>& weights) const;
/// Insert one xz-plane of weights into volume
void insert_weights(const thrust::device_vector<float>& wvec,
unsigned int j,
bool sync,
EDDY::CudaVolume& wvol) const;
/// Elementwise multiplication of weights with predictions.
void make_diagw_p_vectors(const EDDY::CudaVolume& pred,
const thrust::device_vector<float>& wgts,
unsigned int xzp,
bool sync,
thrust::device_vector<float>& wp) const;
/// Premultiply W matrix by diag{w}
void make_diagw_W_matrices(const thrust::device_vector<float>& w,
const thrust::device_vector<float>& W,
unsigned int matsz,
unsigned int nmat,
bool sync,
thrust::device_vector<float>& diagwW) const;
/// Multiply a set of (diag{w}W)' matrices with a bunch of diag{w}p vectors
void make_dwWt_dwp_vectors(const thrust::device_vector<float>& dW,
const thrust::device_vector<float>& dp,
unsigned int matsz,
unsigned int nmat,
bool sync,
thrust::device_vector<float>& dWtdp) const;
/// Make a set of "design matrices" for irregularly sampled splines. GPU.
void make_Wir_matrices(const EDDY::CudaVolume& zcoord,
unsigned int xzp,
bool sync,
thrust::device_vector<float>& Wir) const;
/// Multiply a set of Wir matrices with a set of intensity vectors. GPU.
void make_Wir_t_y_vectors(const EDDY::CudaVolume& y,
const thrust::device_vector<float>& Wir,
unsigned int xzp,
bool sync,
thrust::device_vector<float>& Wirty) const;
/// Multiply a set of Wir matrices by transpose of self and add regularisation matrix. GPU.
void make_WtW_StS_matrices(const thrust::device_vector<float>& Wir,
unsigned int mn,
unsigned int nmat,
const thrust::device_vector<float>& StS,
bool sync,
thrust::device_vector<float>& WtW) const;
/// Solve for spline coefficients
void solve_for_c_hat(// Input
const thrust::device_vector<float>& WtW, // WtW+StS matrices for one xz-plane.
const thrust::device_vector<float>& Wty, // Wty vectors for one xz-plane
unsigned int n, // Size of KtK (nxn)
unsigned int nmat, // Number of matrices for one xz-plane
bool sync, // If true syncs after submitting kernel
// Output
thrust::device_vector<float>& chat) const; // Returns inv(WtW)*Wty for all matrices in xz-plane
/// Multiply spline coefficients with regularly sampled spline matrix
void make_y_hat_vectors(// Input
const thrust::device_vector<float>& W,
const thrust::device_vector<float>& chat,
unsigned int mn,
unsigned int nvec,
bool sync,
// Output
thrust::device_vector<float>& yhat) const;
/// Transfer y_hat (interpolated data) to volume
void transfer_y_hat_to_volume(// Input
const thrust::device_vector<float>& yhat,
unsigned int xzp,
bool sync,
// Output
EDDY::CudaVolume& ovol) const;
/// Sorts with zccord as key and reorders data in the same way
void sort_zcoords_and_intensities(const EDDY::CudaVolume& zcoord,
const EDDY::CudaVolume& data,
bool sync,
thrust::device_vector<float>& szcoord,
thrust::device_vector<float>& sdata) const;
/// Takes sorted vectors of z-coords and values and linearly interpolates onto 0--(zsz-1)
void linear_interpolate_columns(const thrust::device_vector<float>& zcoord,
const thrust::device_vector<float>& val,
unsigned int xsz,
unsigned int ysz,
unsigned int zsz,
bool sync,
thrust::device_vector<float>& ival) const;
void transfer_interpolated_columns_to_volume(const thrust::device_vector<float>& zcols,
bool sync,
EDDY::CudaVolume& vol);
/// i,j->linear index for row-first square matrix. Matrix addressed M(0:mn-1,0:mn-1)
unsigned int rfindx(unsigned int i, unsigned int j, unsigned int mn) const { return(i+j*mn); }
/// i,j->linear index for column-first square matrix. Matrix addressed M(0:mn-1,0:mn-1)
unsigned int cfindx(unsigned int i, unsigned int j, unsigned int mn) const { return(i*mn+j); }
/// Local sqr, just to be safe
template <typename T> T sqr(T a) const { return(a*a); }
/// Writes a single matrix residing on the GPU as a Newmat text file
void write_matrix(const thrust::device_vector<float>& mats,
unsigned int offs,
unsigned int m,
unsigned int n,
const std::string& fname) const;
/// Writes a block of matrices residing on the GPU as a set of Newmat text files
void write_matrices(const thrust::device_vector<float>& mats,
unsigned int nmat,
unsigned int m,
unsigned int n,
const std::string& basefname) const;
/// Writes all neccessary info for one column for debugging
void write_debug_info_for_pred_resampling(unsigned int x,
unsigned int y,
const std::string& bfname,
const EDDY::CudaVolume& z,
const EDDY::CudaVolume& g,
const EDDY::CudaVolume& p,
const thrust::device_vector<float>& sz,
const thrust::device_vector<float>& W,
const thrust::device_vector<float>& Wir,
const thrust::device_vector<float>& w,
const thrust::device_vector<float>& wp,
const thrust::device_vector<float>& wW,
const thrust::device_vector<float>& Wirtg,
const thrust::device_vector<float>& wWtwp,
const thrust::device_vector<float>& WirtWir,
const thrust::device_vector<float>& wWtwW,
const thrust::device_vector<float>& sum_vec,
const thrust::device_vector<float>& sum_mat,
const thrust::device_vector<float>& c_hat,
const thrust::device_vector<float>& y_hat) const;
};
} // End namespace EDDY
cuda/eddy_matrix_kernels_internal.h 0 → 100644
View file @ b9e18ed6
#include "hip/hip_runtime.h"
/*! \file eddy_matrix_kernels_internal.h
\brief Contains declarations of __device__ kernels used internally for LS resampling in Eddy project
The kernels for CUDA based Matrix manipulation have been divided into __global__
ones that have an API from any .c or .cpp file and __device__ ones that are only
callable from within other kernels. The former category are exposed to the application
programmer through EddyMatrixKernels.h and placed in the EddyMatrixKernels namebase.
The latter are "hidden" by virtue of being declared in this .h file which is _only_
intended to be included by EddyMatrixKernels.cu and by residing in the EMKI namespace.
\author Jesper Andersson
\version 1.0b, July, 2013
*/
#ifndef eddy_matrix_kernels_internal_h
#define eddy_matrix_kernels_internal_h
#include <hip/hip_runtime.h>
namespace EMKI {
__device__ void qr_single(// Input
const float *K, // mxn matrix to be QR-decomposed
unsigned int m,
unsigned int n,
// Scratch
float *v, // m elements of scratch space
float *w, // m elements of scratch space
// Thread info
unsigned int id, // Thread id
unsigned int nt, // # of threads for this matrix
// Output
float *Qt, // Q' in K = Q'*R
float *R); // R in K = Q'*R
__device__ void M_eq_M(float *dM, // Destination matrix
const float *sM, // Source matrix
unsigned int m, // No. of rows
unsigned int n, // No. of columns
unsigned int nt, // No. of threads for this matrix
unsigned int id); // Number of this thread (among nt)
__device__ void set_to_identity(float *M, // Matrix to set
unsigned int m, // Size of matrix
unsigned int nt, // Number of threads for this matrix
unsigned int id); // Number of this thread (among nt)
__device__ float get_alfa(const float *M, // Matrix
unsigned int m, // No of rows
unsigned int n, // No of columns
unsigned int j, // Column to get alfa for
unsigned int nt, // No of threads for this matrix
unsigned int id, // Number of thread (among nt)
float *scr);// Scratch space for reduction
__device__ void get_v(const float *M, // The matrix (full size)
unsigned int m, // No. of rows of M
unsigned int n, // No. of columns of M
unsigned int j, // Column to extract v for
float alfa,// Alfa
unsigned int nt, // No. of threads for this matrix
unsigned int id, // Thread number among nt
float *v); // Space for v
__device__ void two_x_vt_x_R(const float *R, // R (first j columns upper diag)
unsigned int m, // No. of rows of R
unsigned int n, // No. of rows of R
unsigned int j, // The column we currently work on
unsigned int nt, // No. of threads for this matrix
unsigned int id, // Thread number (out of nt)
const float *v, // v (first j elements zero)
float *twovtR);// 2.0*v'*R, first j elements zero
__device__ void R_minus_v_x_wt(unsigned int m, // No. of rows of R
unsigned int n, // No. of columns of R
unsigned int j, // The column we are currently on
unsigned int nt, // No. of threads for this matrix
unsigned int id, // Number of this thread (within nt)
const float *v, // v, mx1, j first elements zero
const float *w, // w, nx1, j first elements zero
float *R);// R in, R - vw' out
__device__ void two_x_vt_x_Qt(const float *Qt, // Q'
unsigned int m, // No. of rows of Qt
unsigned int n, // No. of columns of Qt
unsigned int j, // The column we currently work on
unsigned int nt, // No. of threads for this matrix
unsigned int id, // Thread number (out of nt)
const float *v, // v (first j elements zero)
float *twovtQt);// 2.0*v'*Q'
__device__ void Qt_minus_v_x_wt(unsigned int m, // No. of rows of Q'
unsigned int n, // No. of columns of Q'
unsigned int j, // The column we are currently on
unsigned int nt, // No. of threads for this matrix
unsigned int id, // Number of this thread (within nt)
const float *v, // v, mx1, j first elements zero
const float *w, // w, nx1
float *Qt);// Q' in, Q' - vw' out
__device__ bool is_pow_of_two(unsigned int n);
__device__ unsigned int next_pow_of_two(unsigned int n);
__device__ unsigned int rf_indx(unsigned int i, unsigned int j, unsigned int m, unsigned int n);
__device__ unsigned int cf_indx(unsigned int i, unsigned int j, unsigned int m, unsigned int n);
template <typename T>
__device__ T sqr(const T& v) { return(v*v); }
template <typename T>
__device__ T min(const T& p1, const T& p2) { return((p1<p2) ? p1 : p2); }
template <typename T>
__device__ T max(const T& p1, const T& p2) { return((p1<p2) ? p2 : p1); }
__device__ float wgt_at(float x);
__device__ void solve_single(// Input
const float *Qt, // Row-first orthogonal mxm matrix
const float *R, // Row-first upper-diagonal mxn matrix
const float *y, // mx1 vector
unsigned int m,
unsigned int n,
// Output
float *y_hat); // nx1 solution vector
__device__ void back_substitute(const float *R,
unsigned int m,
unsigned int n,
float *v);
__device__ void M_times_v(const float *M,
const float *v,
unsigned int m,
unsigned int n,
float *Mv);
__device__ void cf_KtK_one_mat(const float *K, // Column-first mxn
unsigned int m,
unsigned int n,
const float *StS, // Column-first nxn
float lambda,
unsigned int idr,
unsigned int idc,
unsigned int ntr,
unsigned int ntc,
float *KtK);// Row-first nxn
__device__ void rf_KtK_one_mat(const float *K, // Row-first mxn
unsigned int m,
unsigned int n,
const float *StS, // Row-first nxn
float lambda,
unsigned int idr,
unsigned int idc,
unsigned int ntr,
unsigned int ntc,
float *KtK);// Row-first nxn
__device__ void Ab_one_mat(// Input
const float *A, // Row-first
const float *b,
unsigned int m,
unsigned int n,
unsigned int id,
unsigned int ntr,
// Output
float *Ab);
__device__ void Atb_one_mat(// Input
const float *A, // Row-first
const float *b,
unsigned int m,
unsigned int n,
unsigned int id,
unsigned int ntr,
// Output
float *Atb);
__device__ void Kty_one_mat(// Input
const float *K,
const float *y,
unsigned int m,
unsigned int n,
unsigned int id,
unsigned int ntr,
// Output
float *Kty);
__device__ void Wir_one_mat(// Input
const float *zcoord,
unsigned int zstep,
unsigned int mn,
unsigned int id,
unsigned int ntr,
// Output
float *Wir);
__device__ void Wirty_one_mat(// Input
const float *y,
const float *Wir, // Row-first
unsigned int zstep,
unsigned int mn,
unsigned int id,
unsigned int ntr,
// Output
float *Wirty);
} // End namespace EMKI
#endif // End #ifndef eddy_matrix_kernels_internal_h
cuda/eddy_matrix_kernels_internal.h.prehip 0 → 100644
View file @ b9e18ed6
/*! \file eddy_matrix_kernels_internal.h
\brief Contains declarations of __device__ kernels used internally for LS resampling in Eddy project
The kernels for CUDA based Matrix manipulation have been divided into __global__
ones that have an API from any .c or .cpp file and __device__ ones that are only
callable from within other kernels. The former category are exposed to the application
programmer through EddyMatrixKernels.h and placed in the EddyMatrixKernels namebase.
The latter are "hidden" by virtue of being declared in this .h file which is _only_
intended to be included by EddyMatrixKernels.cu and by residing in the EMKI namespace.
\author Jesper Andersson
\version 1.0b, July, 2013
*/
#ifndef eddy_matrix_kernels_internal_h
#define eddy_matrix_kernels_internal_h
#include <cuda.h>
namespace EMKI {
__device__ void qr_single(// Input
const float *K, // mxn matrix to be QR-decomposed
unsigned int m,
unsigned int n,
// Scratch
float *v, // m elements of scratch space
float *w, // m elements of scratch space
// Thread info
unsigned int id, // Thread id
unsigned int nt, // # of threads for this matrix
// Output
float *Qt, // Q' in K = Q'*R
float *R); // R in K = Q'*R
__device__ void M_eq_M(float *dM, // Destination matrix
const float *sM, // Source matrix
unsigned int m, // No. of rows
unsigned int n, // No. of columns
unsigned int nt, // No. of threads for this matrix
unsigned int id); // Number of this thread (among nt)
__device__ void set_to_identity(float *M, // Matrix to set
unsigned int m, // Size of matrix
unsigned int nt, // Number of threads for this matrix
unsigned int id); // Number of this thread (among nt)
__device__ float get_alfa(const float *M, // Matrix
unsigned int m, // No of rows
unsigned int n, // No of columns
unsigned int j, // Column to get alfa for
unsigned int nt, // No of threads for this matrix
unsigned int id, // Number of thread (among nt)
float *scr);// Scratch space for reduction
__device__ void get_v(const float *M, // The matrix (full size)
unsigned int m, // No. of rows of M
unsigned int n, // No. of columns of M
unsigned int j, // Column to extract v for
float alfa,// Alfa
unsigned int nt, // No. of threads for this matrix
unsigned int id, // Thread number among nt
float *v); // Space for v
__device__ void two_x_vt_x_R(const float *R, // R (first j columns upper diag)
unsigned int m, // No. of rows of R
unsigned int n, // No. of rows of R
unsigned int j, // The column we currently work on
unsigned int nt, // No. of threads for this matrix
unsigned int id, // Thread number (out of nt)
const float *v, // v (first j elements zero)
float *twovtR);// 2.0*v'*R, first j elements zero
__device__ void R_minus_v_x_wt(unsigned int m, // No. of rows of R
unsigned int n, // No. of columns of R
unsigned int j, // The column we are currently on
unsigned int nt, // No. of threads for this matrix
unsigned int id, // Number of this thread (within nt)
const float *v, // v, mx1, j first elements zero
const float *w, // w, nx1, j first elements zero
float *R);// R in, R - vw' out
__device__ void two_x_vt_x_Qt(const float *Qt, // Q'
unsigned int m, // No. of rows of Qt
unsigned int n, // No. of columns of Qt
unsigned int j, // The column we currently work on
unsigned int nt, // No. of threads for this matrix
unsigned int id, // Thread number (out of nt)
const float *v, // v (first j elements zero)
float *twovtQt);// 2.0*v'*Q'
__device__ void Qt_minus_v_x_wt(unsigned int m, // No. of rows of Q'
unsigned int n, // No. of columns of Q'
unsigned int j, // The column we are currently on
unsigned int nt, // No. of threads for this matrix
unsigned int id, // Number of this thread (within nt)
const float *v, // v, mx1, j first elements zero
const float *w, // w, nx1
float *Qt);// Q' in, Q' - vw' out
__device__ bool is_pow_of_two(unsigned int n);
__device__ unsigned int next_pow_of_two(unsigned int n);
__device__ unsigned int rf_indx(unsigned int i, unsigned int j, unsigned int m, unsigned int n);
__device__ unsigned int cf_indx(unsigned int i, unsigned int j, unsigned int m, unsigned int n);
template <typename T>
__device__ T sqr(const T& v) { return(v*v); }
template <typename T>
__device__ T min(const T& p1, const T& p2) { return((p1<p2) ? p1 : p2); }
template <typename T>
__device__ T max(const T& p1, const T& p2) { return((p1<p2) ? p2 : p1); }
__device__ float wgt_at(float x);
__device__ void solve_single(// Input
const float *Qt, // Row-first orthogonal mxm matrix
const float *R, // Row-first upper-diagonal mxn matrix
const float *y, // mx1 vector
unsigned int m,
unsigned int n,
// Output
float *y_hat); // nx1 solution vector
__device__ void back_substitute(const float *R,
unsigned int m,
unsigned int n,
float *v);
__device__ void M_times_v(const float *M,
const float *v,
unsigned int m,
unsigned int n,
float *Mv);
__device__ void cf_KtK_one_mat(const float *K, // Column-first mxn
unsigned int m,
unsigned int n,
const float *StS, // Column-first nxn
float lambda,
unsigned int idr,
unsigned int idc,
unsigned int ntr,
unsigned int ntc,
float *KtK);// Row-first nxn
__device__ void rf_KtK_one_mat(const float *K, // Row-first mxn
unsigned int m,
unsigned int n,
const float *StS, // Row-first nxn
float lambda,
unsigned int idr,
unsigned int idc,
unsigned int ntr,
unsigned int ntc,
float *KtK);// Row-first nxn
__device__ void Ab_one_mat(// Input
const float *A, // Row-first
const float *b,
unsigned int m,
unsigned int n,
unsigned int id,
unsigned int ntr,
// Output
float *Ab);
__device__ void Atb_one_mat(// Input
const float *A, // Row-first
const float *b,
unsigned int m,
unsigned int n,
unsigned int id,
unsigned int ntr,
// Output
float *Atb);
__device__ void Kty_one_mat(// Input
const float *K,
const float *y,
unsigned int m,
unsigned int n,
unsigned int id,
unsigned int ntr,
// Output
float *Kty);
__device__ void Wir_one_mat(// Input
const float *zcoord,
unsigned int zstep,
unsigned int mn,
unsigned int id,
unsigned int ntr,
// Output
float *Wir);
__device__ void Wirty_one_mat(// Input
const float *y,
const float *Wir, // Row-first
unsigned int zstep,
unsigned int mn,
unsigned int id,
unsigned int ntr,
// Output
float *Wirty);
} // End namespace EMKI
#endif // End #ifndef eddy_matrix_kernels_internal_h
cuda/fmriPredictor.cpp 0 → 100644
View file @ b9e18ed6
/*! \file fmriPredictor.cu
\brief Contains definitions for class for making Gaussian process based predictions about fMRI data.
\author Jesper Andersson
\version 1.0b, May., 2022.
*/
// Definitions of class to make Gaussian-Process
// based predictions about fMRI data.
//
// fmriPredictor.cu
//
// Jesper Andersson, FMRIB Image Analysis Group
//
// Copyright (C) 2011 University of Oxford
//
#include <cstdlib>
#include <string>
#include <vector>
#include <cmath>
#pragma push
#pragma diag_suppress = code_is_unreachable // Supress warnings from armawrap
#pragma diag_suppress = expr_has_no_effect // Supress warnings from boost
#include "armawrap/newmat.h"
#include "newimage/newimageall.h"
#pragma pop
#include "miscmaths/miscmaths.h"
#include "EddyHelperClasses.h"
#include "EddyUtils.h"
#include "fmriPredictor.h"
#include "CudaVolume.h"
using namespace EDDY;
/****************************************************************//**
* \brief Returns prediction for point given by indx
*
* Returns a predicted image for the given index, where index refers
* to the corresponding time-point.
* It should be noted that this implementation isn't very efficient as
* it reloads all image to the GPU. That means that if there are N images
* and we want to predict them all it takes N*N transfers to the GPU.
* \param index refers to the corresponding time-point
* \param exclude Decides if indx itself should be used in the prediction
* (exclude=false) or not (exclude=true)
* \param pvec Prediction vector for index
* \param pi The "predicted image"
*
********************************************************************/
void fmriPredictor::predict_image_gpu(// Input
unsigned int indx,
bool exclude,
const arma::rowvec& pvec,
// Output
NEWIMAGE::volume<float>& pi) const EddyTry
{
unsigned int pvec_length = (exclude) ? this->no_of_scans_in_same_session(indx)-1 : this->no_of_scans_in_same_session(indx);
if (pvec_length != pvec.n_cols) throw EddyException("fmriPredictor::predict_image_gpu: Wrong length pvec");
if (!NEWIMAGE::samesize(pi,*(this->first_imptr()))) {
pi.reinitialize(this->first_imptr()->xsize(),this->first_imptr()->ysize(),this->first_imptr()->zsize());
NEWIMAGE::copybasicproperties(*(this->first_imptr()),pi);
}
EDDY::CudaVolume pcv(pi,false);
int sindx = this->index_in_session(indx);
std::shared_ptr<NEWIMAGE::volume<float> > ptr=nullptr;
for (unsigned int s=0; s<this->no_of_scans_in_same_session(indx); s++) {
if (exclude) {
if (s < sindx) pcv.MultiplyAndAddToMe(EDDY::CudaVolume(*(this->imptr_from_same_session(indx,s))),pvec(s));
// Do nothing if s==sindx
else if (s > indx) pcv.MultiplyAndAddToMe(EDDY::CudaVolume(*(this->imptr_from_same_session(indx,s))),pvec(s-1));
}
else pcv.MultiplyAndAddToMe(EDDY::CudaVolume(*(this->imptr_from_same_session(indx,s))),pvec(s));
}
pcv += EDDY::CudaVolume(*(this->meanptr_from_same_session(indx)));
pcv.GetVolume(pi);
return;
} EddyCatch
void fmriPredictor::predict_images_gpu(// Input
const std::vector<unsigned int>& indicies,
bool exclude,
const std::vector<arma::rowvec>& pvecs,
// Output
std::vector<NEWIMAGE::volume<float> >& pi) const EddyTry
{
// Sanity checks
if (indicies.size() != pvecs.size() || indicies.size() != pi.size()) {
throw EDDY::EddyException("fmriPredictor::predict_images_gpu: mismatch among indicies, pvecs and pi");
}
if (!this->all_in_same_session(indicies)) throw EDDY::EddyException("fmriPredictor::predict_images_gpu: indicies are in different sessions");
unsigned int pvec_length = (exclude) ? this->no_of_scans_in_same_session(indicies[0])-1 : this->no_of_scans_in_same_session(indicies[0]);
if (std::any_of(pvecs.begin(),pvecs.end(),[pvec_length](const arma::rowvec& rv){ return(rv.n_cols != pvec_length); })) {
throw EDDY::EddyException("fmriPredictor::predict_images_gpu: pvec with wrong length");
}
// Start by allocating space on GPU for all output images
std::vector<EDDY::CudaVolume> pcvs(indicies.size());
for (unsigned int i=0; i<indicies.size(); i++) {
if (!NEWIMAGE::samesize(pi[i],*(this->first_imptr()))) {
pi[i].reinitialize(this->first_imptr()->xsize(),this->first_imptr()->ysize(),this->first_imptr()->zsize());
NEWIMAGE::copybasicproperties(*(this->first_imptr()),pi[i]);
}
pcvs[i].SetHdr(pi[i]);
}
// Transfer mean image to the GPU
EDDY::CudaVolume mean = *(this->meanptr_from_same_session(indicies[0]));
// Do the GP predictions
for (unsigned int s=0; s<this->no_of_scans_in_same_session(indicies[0]); s++) { // s index into original volumes
EDDY::CudaVolume cv = *(this->imptr_from_same_session(indicies[0],s));
for (unsigned int i=0; i<indicies.size(); i++) { // i index into predictions
if (exclude) {
if (s < this->index_in_session(indicies[i])) pcvs[i].MultiplyAndAddToMe(cv,pvecs[i](s));
// Do nothing if (s == indicies[i])
else if (s > this->index_in_session(indicies[i])) pcvs[i].MultiplyAndAddToMe(cv,pvecs[i](s-1));
}
else pcvs[i].MultiplyAndAddToMe(cv,pvecs[i](s));
}
}
// Add means to predictions and transfer back from GPU
for (unsigned int i=0; i<indicies.size(); i++) {
pcvs[i] += mean;
pcvs[i].GetVolume(pi[i]);
}
return;
} EddyCatch
cuda/fmriPredictor.cu.prehip 0 → 100644
View file @ b9e18ed6
/*! \file fmriPredictor.cu
\brief Contains definitions for class for making Gaussian process based predictions about fMRI data.
\author Jesper Andersson
\version 1.0b, May., 2022.
*/
// Definitions of class to make Gaussian-Process
// based predictions about fMRI data.
//
// fmriPredictor.cu
//
// Jesper Andersson, FMRIB Image Analysis Group
//
// Copyright (C) 2011 University of Oxford
//
#include <cstdlib>
#include <string>
#include <vector>
#include <cmath>
#pragma push
#pragma diag_suppress = code_is_unreachable // Supress warnings from armawrap
#pragma diag_suppress = expr_has_no_effect // Supress warnings from boost
#include "armawrap/newmat.h"
#include "newimage/newimageall.h"
#pragma pop
#include "miscmaths/miscmaths.h"
#include "EddyHelperClasses.h"
#include "EddyUtils.h"
#include "fmriPredictor.h"
#include "CudaVolume.h"
using namespace EDDY;
/****************************************************************//**
* \brief Returns prediction for point given by indx
*
* Returns a predicted image for the given index, where index refers
* to the corresponding time-point.
* It should be noted that this implementation isn't very efficient as
* it reloads all image to the GPU. That means that if there are N images
* and we want to predict them all it takes N*N transfers to the GPU.
* \param index refers to the corresponding time-point
* \param exclude Decides if indx itself should be used in the prediction
* (exclude=false) or not (exclude=true)
* \param pvec Prediction vector for index
* \param pi The "predicted image"
*
********************************************************************/
void fmriPredictor::predict_image_gpu(// Input
unsigned int indx,
bool exclude,
const arma::rowvec& pvec,
// Output
NEWIMAGE::volume<float>& pi) const EddyTry
{
unsigned int pvec_length = (exclude) ? this->no_of_scans_in_same_session(indx)-1 : this->no_of_scans_in_same_session(indx);
if (pvec_length != pvec.n_cols) throw EddyException("fmriPredictor::predict_image_gpu: Wrong length pvec");
if (!NEWIMAGE::samesize(pi,*(this->first_imptr()))) {
pi.reinitialize(this->first_imptr()->xsize(),this->first_imptr()->ysize(),this->first_imptr()->zsize());
NEWIMAGE::copybasicproperties(*(this->first_imptr()),pi);
}
EDDY::CudaVolume pcv(pi,false);
int sindx = this->index_in_session(indx);
std::shared_ptr<NEWIMAGE::volume<float> > ptr=nullptr;
for (unsigned int s=0; s<this->no_of_scans_in_same_session(indx); s++) {
if (exclude) {
if (s < sindx) pcv.MultiplyAndAddToMe(EDDY::CudaVolume(*(this->imptr_from_same_session(indx,s))),pvec(s));
// Do nothing if s==sindx
else if (s > indx) pcv.MultiplyAndAddToMe(EDDY::CudaVolume(*(this->imptr_from_same_session(indx,s))),pvec(s-1));
}
else pcv.MultiplyAndAddToMe(EDDY::CudaVolume(*(this->imptr_from_same_session(indx,s))),pvec(s));
}
pcv += EDDY::CudaVolume(*(this->meanptr_from_same_session(indx)));
pcv.GetVolume(pi);
return;
} EddyCatch
void fmriPredictor::predict_images_gpu(// Input
const std::vector<unsigned int>& indicies,
bool exclude,
const std::vector<arma::rowvec>& pvecs,
// Output
std::vector<NEWIMAGE::volume<float> >& pi) const EddyTry
{
// Sanity checks
if (indicies.size() != pvecs.size() || indicies.size() != pi.size()) {
throw EDDY::EddyException("fmriPredictor::predict_images_gpu: mismatch among indicies, pvecs and pi");
}
if (!this->all_in_same_session(indicies)) throw EDDY::EddyException("fmriPredictor::predict_images_gpu: indicies are in different sessions");
unsigned int pvec_length = (exclude) ? this->no_of_scans_in_same_session(indicies[0])-1 : this->no_of_scans_in_same_session(indicies[0]);
if (std::any_of(pvecs.begin(),pvecs.end(),[pvec_length](const arma::rowvec& rv){ return(rv.n_cols != pvec_length); })) {
throw EDDY::EddyException("fmriPredictor::predict_images_gpu: pvec with wrong length");
}
// Start by allocating space on GPU for all output images
std::vector<EDDY::CudaVolume> pcvs(indicies.size());
for (unsigned int i=0; i<indicies.size(); i++) {
if (!NEWIMAGE::samesize(pi[i],*(this->first_imptr()))) {
pi[i].reinitialize(this->first_imptr()->xsize(),this->first_imptr()->ysize(),this->first_imptr()->zsize());
NEWIMAGE::copybasicproperties(*(this->first_imptr()),pi[i]);
}
pcvs[i].SetHdr(pi[i]);
}
// Transfer mean image to the GPU
EDDY::CudaVolume mean = *(this->meanptr_from_same_session(indicies[0]));
// Do the GP predictions
for (unsigned int s=0; s<this->no_of_scans_in_same_session(indicies[0]); s++) { // s index into original volumes
EDDY::CudaVolume cv = *(this->imptr_from_same_session(indicies[0],s));
for (unsigned int i=0; i<indicies.size(); i++) { // i index into predictions
if (exclude) {
if (s < this->index_in_session(indicies[i])) pcvs[i].MultiplyAndAddToMe(cv,pvecs[i](s));
// Do nothing if (s == indicies[i])
else if (s > this->index_in_session(indicies[i])) pcvs[i].MultiplyAndAddToMe(cv,pvecs[i](s-1));
}
else pcvs[i].MultiplyAndAddToMe(cv,pvecs[i](s));
}
}
// Add means to predictions and transfer back from GPU
for (unsigned int i=0; i<indicies.size(); i++) {
pcvs[i] += mean;
pcvs[i].GetVolume(pi[i]);
}
return;
} EddyCatch
cuda/math_constants.h 0 → 100644
View file @ b9e18ed6
#if !defined(__MATH_CONSTANTS_H__)
#define __MATH_CONSTANTS_H__
#define CUDART_INF_F __int_as_float(0x7f800000U)
#define CUDART_NAN_F __int_as_float(0x7fffffffU)
#define CUDART_MIN_DENORM_F __int_as_float(0x00000001U)
#define CUDART_MAX_NORMAL_F __int_as_float(0x7f7fffffU)
#define CUDART_NEG_ZERO_F __int_as_float(0x80000000U)
#define CUDART_ZERO_F 0.0F
#define CUDART_ONE_F 1.0F
#define CUDART_SQRT_HALF_F 0.707106781F
#define CUDART_SQRT_HALF_HI_F 0.707106781F
#define CUDART_SQRT_HALF_LO_F 1.210161749e-08F
#define CUDART_SQRT_TWO_F 1.414213562F
#define CUDART_THIRD_F 0.333333333F
#define CUDART_PIO4_F 0.785398163F
#define CUDART_PIO2_F 1.570796327F
#define CUDART_3PIO4_F 2.356194490F
#define CUDART_2_OVER_PI_F 0.636619772F
#define CUDART_SQRT_2_OVER_PI_F 0.797884561F
#define CUDART_PI_F 3.141592654F
#define CUDART_L2E_F 1.442695041F
#define CUDART_L2T_F 3.321928094F
#define CUDART_LG2_F 0.301029996F
#define CUDART_LGE_F 0.434294482F
#define CUDART_LN2_F 0.693147181F
#define CUDART_LNT_F 2.302585093F
#define CUDART_LNPI_F 1.144729886F
#define CUDART_TWO_TO_M126_F 1.175494351e-38F
#define CUDART_TWO_TO_126_F 8.507059173e37F
#define CUDART_NORM_HUGE_F 3.402823466e38F
#define CUDART_TWO_TO_23_F 8388608.0F
#define CUDART_TWO_TO_24_F 16777216.0F
#define CUDART_TWO_TO_31_F 2147483648.0F
#define CUDART_TWO_TO_32_F 4294967296.0F
#define CUDART_REMQUO_BITS_F 3U
#define CUDART_REMQUO_MASK_F (~((~0U)<<CUDART_REMQUO_BITS_F))
#define CUDART_TRIG_PLOSS_F 105615.0F
#define CUDART_INF __longlong_as_double(0x7ff0000000000000ULL)
#define CUDART_NAN __longlong_as_double(0xfff8000000000000ULL)
#define CUDART_NEG_ZERO __longlong_as_double(0x8000000000000000ULL)
#define CUDART_MIN_DENORM __longlong_as_double(0x0000000000000001ULL)
#define CUDART_ZERO 0.0
#define CUDART_ONE 1.0
#define CUDART_SQRT_TWO 1.4142135623730951e+0
#define CUDART_SQRT_HALF 7.0710678118654757e-1
#define CUDART_SQRT_HALF_HI 7.0710678118654757e-1
#define CUDART_SQRT_HALF_LO (-4.8336466567264567e-17)
#define CUDART_THIRD 3.3333333333333333e-1
#define CUDART_TWOTHIRD 6.6666666666666667e-1
#define CUDART_PIO4 7.8539816339744828e-1
#define CUDART_PIO4_HI 7.8539816339744828e-1
#define CUDART_PIO4_LO 3.0616169978683830e-17
#define CUDART_PIO2 1.5707963267948966e+0
#define CUDART_PIO2_HI 1.5707963267948966e+0
#define CUDART_PIO2_LO 6.1232339957367660e-17
#define CUDART_3PIO4 2.3561944901923448e+0
#define CUDART_2_OVER_PI 6.3661977236758138e-1
#define CUDART_PI 3.1415926535897931e+0
#define CUDART_PI_HI 3.1415926535897931e+0
#define CUDART_PI_LO 1.2246467991473532e-16
#define CUDART_SQRT_2PI 2.5066282746310007e+0
#define CUDART_SQRT_2PI_HI 2.5066282746310007e+0
#define CUDART_SQRT_2PI_LO (-1.8328579980459167e-16)
#define CUDART_SQRT_PIO2 1.2533141373155003e+0
#define CUDART_SQRT_PIO2_HI 1.2533141373155003e+0
#define CUDART_SQRT_PIO2_LO (-9.1642899902295834e-17)
#define CUDART_SQRT_2OPI 7.9788456080286536e-1
#define CUDART_L2E 1.4426950408889634e+0
#define CUDART_L2E_HI 1.4426950408889634e+0
#define CUDART_L2E_LO 2.0355273740931033e-17
#define CUDART_L2T 3.3219280948873622e+0
#define CUDART_LG2 3.0102999566398120e-1
#define CUDART_LG2_HI 3.0102999566398120e-1
#define CUDART_LG2_LO (-2.8037281277851704e-18)
#define CUDART_LGE 4.3429448190325182e-1
#define CUDART_LGE_HI 4.3429448190325182e-1
#define CUDART_LGE_LO 1.09831965021676510e-17
#define CUDART_LN2 6.9314718055994529e-1
#define CUDART_LN2_HI 6.9314718055994529e-1
#define CUDART_LN2_LO 2.3190468138462996e-17
#define CUDART_LNT 2.3025850929940459e+0
#define CUDART_LNT_HI 2.3025850929940459e+0
#define CUDART_LNT_LO (-2.1707562233822494e-16)
#define CUDART_LNPI 1.1447298858494002e+0
#define CUDART_LN2_X_1024 7.0978271289338397e+2
#define CUDART_LN2_X_1025 7.1047586007394398e+2
#define CUDART_LN2_X_1075 7.4513321910194122e+2
#define CUDART_LG2_X_1024 3.0825471555991675e+2
#define CUDART_LG2_X_1075 3.2360724533877976e+2
#define CUDART_TWO_TO_23 8388608.0
#define CUDART_TWO_TO_52 4503599627370496.0
#define CUDART_TWO_TO_53 9007199254740992.0
#define CUDART_TWO_TO_54 18014398509481984.0
#define CUDART_TWO_TO_M54 5.5511151231257827e-17
#define CUDART_TWO_TO_M1022 2.22507385850720140e-308
#define CUDART_TRIG_PLOSS 2147483648.0
#define CUDART_DBL2INT_CVT 6755399441055744.0
#endif
eddy 0 → 100755
View file @ b9e18ed6
#!/usr/bin/env fslpython
#
# Call an appropriate variant of eddy based on
# what is installed, and whether we have a GPU.
import sys
import os
import os.path as op
import subprocess as sp
def main():
"""Call an appropriate variant of eddy based on what is installed, and
whether a CUDA driver is available.
"""
fsldir = os.environ['FSLDIR']
find_cuda = op.join(fsldir, 'bin', 'find_cuda_exe')
eddy_cuda = sp.check_output((find_cuda, 'eddy_cuda'), text=True).strip()
eddy_cpu = op.join(fsldir, 'bin', 'eddy_cpu')
# Call eddy_cuda if it exists and we
# are running on a system with a GPU
if eddy_cuda != '':
cmd = eddy_cuda
# Otherwise fall back to eddy_cpu
elif op.exists(eddy_cpu):
cmd = eddy_cpu
else:
print(f'Cannot find suitable eddy executable in {fsldir}/bin/!')
sys.exit(1)
os.execl(cmd, cmd, *sys.argv[1:])
if __name__ == '__main__':
main()
eddy.cpp 0 → 100644
View file @ b9e18ed6
/*! \file eddy.cpp
\brief Contains main() and some very high level functions for eddy
*/
#if __cplusplus >= 201103L || __clang__
#include <array>
using std::array;
#else
#include <tr1/array>
using std::tr1::array;
#endif
#include <cstdlib>
#include <string>
#include <vector>
#include <cmath>
#include <memory>
#include <thread>
#include "nlohmann/json.hpp"
#include "armawrap/newmat.h"
#include "newimage/newimageall.h"
#include "miscmaths/miscmaths.h"
#include "utils/stack_dump.h"
#include "utils/FSLProfiler.h"
#include "EddyHelperClasses.h"
#include "ECScanClasses.h"
#include "DiffusionGP.h"
#include "b0Predictor.h"
#include "fmriPredictor.h"
#include "EddyUtils.h"
#include "EddyCommandLineOptions.h"
#include "PostEddyCF.h"
#include "PostEddyAlignShellsFunctions.h"
#include "MoveBySuscCF.h"
#include "BiasFieldEstimator.h"
#include "eddy.h"
#ifdef COMPILE_GPU
#include "cuda/GpuPredictorChunk.h"
#include "cuda/EddyGpuUtils.h"
#endif
namespace EDDY {
/// A struct/class that gathers nlohmann::json objects from the various file writing functions.
/// These are then written out as a single .json file. Longer term the plan is to move away
/// from writing individual text files in favour of a single .json file. It is declared here
/// as part of hiding nlohmann from nvcc.
class json_out_struct
{
public:
nlohmann::json shell_indicies;
nlohmann::json outlier_report;
std::vector<nlohmann::json> outlier_maps;
nlohmann::json parameters;
nlohmann::json movement_over_time;
nlohmann::json movement_rms;
nlohmann::json restricted_movement_rms;
nlohmann::json long_ec_parameters;
void Write(const std::string& fname) const;
};
} // End namespace EDDY
using namespace std;
using namespace EDDY;
/// The entry point of eddy.
int main(int argc, char *argv[]) try
{
StackDump::Install(); // Gives us informative stack dump if/when program crashes
// Parse comand line input
EddyCommandLineOptions clo(argc,argv); // Command Line Options
// Set range of b-values considered to be in the same shell
if (clo.IsDiffusion()) EddyUtils::SetbRange(clo.BValueRange());
// Prime profiler if requested by user
if (clo.LogTimings()) Utilities::FSLProfiler::SetProfilingOn(clo.LoggerFname());
Utilities::FSLProfiler prof("_"+string(__FILE__)+"_"+string(__func__));
double total_key = prof.StartEntry("Begins eddy run");
EDDY::json_out_struct json_out; // Structure for gathering json output objects.
// Read all available info
if (clo.Verbose()) cout << "Reading images" << endl;
ECScanManager sm = (clo.IsDiffusion()) ? // Conditional assignment because ECScanManager has no default constructor
ECScanManager(clo.ImaFname(),clo.MaskFname(),clo.AcqpFname(),clo.TopupFname(),
clo.FieldFname(),clo.FieldMatFname(),clo.BVecsFname(),clo.BValsFname(),
clo.BDeltasFname(),clo.RepetitionTime(),clo.EchoTime(),clo.FirstLevelModel(),
clo.b0_FirstLevelModel(),clo.LongECModel(),clo.Indicies(),clo.SessionIndicies(),
clo.PolationParameters(),clo.MultiBand(),clo.DontCheckShelling()) // Scan Manager for diffusion
: // If clo.IsDiffusion() is false
ECScanManager(clo.ImaFname(),clo.MaskFname(),clo.AcqpFname(),clo.TopupFname(),
clo.FieldFname(),clo.FieldMatFname(),clo.RepetitionTime(),
clo.EchoTime(),clo.Indicies(),clo.SessionIndicies(),
clo.PolationParameters(),clo.MultiBand()); // Scan Manager for fmri
if (clo.FillEmptyPlanes()) { if (clo.Verbose()) cout << "Filling empty planes" << endl; sm.FillEmptyPlanes(); }
if (clo.ResamplingMethod() == FinalResamplingType::LSR) {
if (!sm.CanDoLSRResampling()) throw EddyException("These data do not support least-squares resampling");
}
if (clo.UseB0sToAlignShellsPostEddy() && !sm.B0sAreUsefulForPEAS()) {
throw EddyException("These data do not support using b0s for Post Eddy Alignment of Shells");
}
if (clo.RefScanNumber()) sm.SetLocationReference(clo.RefScanNumber());
// Write .json file with shell indicies for use by EddyQC and debugging
if (clo.IsDiffusion()) json_out.shell_indicies = sm.WriteShellIndicies(clo.ShellIndexOutFname());
// Write topup-field if debug flag is set
if (clo.DebugLevel() && sm.HasSuscHzOffResField()) {
std::string fname = "EDDY_DEBUG_susc_00_0000"; NEWIMAGE::write_volume(*(sm.GetSuscHzOffResField()),fname);
}
// Set initial parameters. This option is only for testing/debugging/personal use
if (clo.InitFname() != std::string("")) {
if (clo.RegisterDWI() && clo.Registerb0()) sm.SetParameters(clo.InitFname(),ScanType::Any);
else if (clo.RegisterDWI()) sm.SetParameters(clo.InitFname(),ScanType::DWI);
else sm.SetParameters(clo.InitFname(),ScanType::b0);
}
// Do the registration
//////////////////////////////////////////////////////////////////////
// The first, and possibly only, registration step is volume_to_volume
//////////////////////////////////////////////////////////////////////
double vol_key = prof.StartEntry("Calling DoVolumeToVolumeRegistration");
if (clo.Verbose()) cout << "Performing volume-to-volume registration" << endl;
ReplacementManager *dwi_rm=NULL;
if (clo.EstimateMoveBySusc()) { // Restrict EC if we are to eventually estimate MBS
EDDY::SecondLevelECModelType b0_slm = clo.b0_SecondLevelModel();
EDDY::SecondLevelECModelType dwi_slm = clo.SecondLevelModel();
if (clo.VeryVerbose()) cout << "Setting linear second level model" << endl;
clo.Set_b0_SecondLevelModel(SecondLevelECModelType::Linear);
clo.SetSecondLevelModel(SecondLevelECModelType::Linear);
dwi_rm = DoVolumeToVolumeRegistration(clo,sm);
if (clo.VeryVerbose()) cout << "Resetting second level model" << endl;
clo.Set_b0_SecondLevelModel(b0_slm);
clo.SetSecondLevelModel(dwi_slm);
}
else dwi_rm = DoVolumeToVolumeRegistration(clo,sm);
prof.EndEntry(vol_key);
// The remaining steps we only run if registration was actually run,
// i.e. not if we read an initialisation file and did zero iterations.
if (clo.NIter() > 0) {
sm.ApplyLocationReference();
// Write text-file with MI values if requested (testing/debugging)
if (clo.PrintMIValues()) {
if (clo.Verbose()) cout << "Writing MI values between shells" << endl;
PrintMIValues(clo,sm,clo.MIPrintFname(),clo.PrintMIPlanes());
}
// Check for residual position differences between shells
if (clo.RegisterDWI()) {
double peas_key = prof.StartEntry("Calling PostEddyAlignShells");
if (!clo.SeparateOffsetFromMovement() && !clo.AlignShellsPostEddy()) {
if (clo.Verbose()) cout << "Checking shell alignment along PE-direction (running PostEddyAlignShellsAlongPE)" << endl;
PEASUtils::PostEddyAlignShellsAlongPE(clo,false,sm);
if (clo.Verbose()) cout << "Checking shell alignment (running PostEddyAlignShells)" << endl;
PEASUtils::PostEddyAlignShells(clo,false,sm);
}
else if (clo.SeparateOffsetFromMovement() && !clo.AlignShellsPostEddy()) {
if (clo.Verbose()) cout << "Aligning shells along PE-direction (running PostEddyAlignShellsAlongPE)" << endl;
PEASUtils::PostEddyAlignShellsAlongPE(clo,true,sm);
if (clo.Verbose()) cout << "Checking shell alignment (running PostEddyAlignShells)" << endl;
PEASUtils::PostEddyAlignShells(clo,false,sm);
}
else if (clo.AlignShellsPostEddy()) {
if (clo.Verbose()) cout << "Checking shell alignment along PE-direction (running PostEddyAlignShellsAlongPE)" << endl;
PEASUtils::PostEddyAlignShellsAlongPE(clo,false,sm);
if (clo.Verbose()) cout << "Aligning shells (running PostEddyAlignShells)" << endl;
PEASUtils::PostEddyAlignShells(clo,true,sm);
}
prof.EndEntry(peas_key);
}
}
//////////////////////////////////////////////////////////////////////
// Next, estimate long time-constant EC if requested
//////////////////////////////////////////////////////////////////////
bool original_sep_offs_move = clo.SeparateOffsetFromMovement();
if (clo.EstimateLongEC()) {
double longEC_key = prof.StartEntry("Calling EstimateLongEC");
if (clo.Verbose()) {
cout << "Estimating long time-constant eddy currents" << endl;
cout << "Using model: " << clo.LongECModelString() << endl;
if (clo.ReestimateECWhenEstimatingLongEC()) {
if (clo.Verbose()) cout << "Also jointly re-estimating EC and movement parameters" << endl;
}
}
// The long-EC estimates _should_ break the very strong covariance between estimates
// of EC-DC-component and subject movement along PE-direction. There is therefore an
// option to "turn off" the attempted separation of EC-DC and subject movement at this
// point. N.B. that it will remain turned off after this point. But we will still align
// different shells along PE. Messy and ugly, and should be refactored at some point.
// I leave it like this for now to avoid changing the interface (breaking scripts).
if (!clo.SeparateOffsetFromMovementWhenEstimatingLongEC()) {
clo.SetSeparateOffsetFromMovement(false);
}
dwi_rm = EstimateLongEC(clo,sm,dwi_rm);
prof.EndEntry(longEC_key);
sm.ApplyLocationReference();
// Check for residual position differences between shells
if (clo.RegisterDWI()) {
double peas_key = prof.StartEntry("Calling PostEddyAlignShells");
if (!original_sep_offs_move && !clo.AlignShellsPostEddy()) {
if (clo.Verbose()) cout << "Checking shell alignment along PE-direction (running PostEddyAlignShellsAlongPE)" << endl;
PEASUtils::PostEddyAlignShellsAlongPE(clo,false,sm);
if (clo.Verbose()) cout << "Checking shell alignment (running PostEddyAlignShells)" << endl;
PEASUtils::PostEddyAlignShells(clo,false,sm);
}
else if (original_sep_offs_move && !clo.AlignShellsPostEddy()) {
if (clo.Verbose()) cout << "Aligning shells along PE-direction (running PostEddyAlignShellsAlongPE)" << endl;
PEASUtils::PostEddyAlignShellsAlongPE(clo,true,sm);
if (clo.Verbose()) cout << "Checking shell alignment (running PostEddyAlignShells)" << endl;
PEASUtils::PostEddyAlignShells(clo,false,sm);
}
else if (clo.AlignShellsPostEddy()) {
if (clo.Verbose()) cout << "Checking shell alignment along PE-direction (running PostEddyAlignShellsAlongPE)" << endl;
PEASUtils::PostEddyAlignShellsAlongPE(clo,false,sm);
if (clo.Verbose()) cout << "Aligning shells (running PostEddyAlignShells)" << endl;
PEASUtils::PostEddyAlignShells(clo,true,sm);
}
prof.EndEntry(peas_key);
}
}
//////////////////////////////////////////////////////////////////////
// Next do a first round of estimating a bias field if requested
//////////////////////////////////////////////////////////////////////
// EstimateBiasField(clo,10.0,1.0,sm);
// std::shared_ptr<const NEWIMAGE::volume<float> > bfield = sm.GetBiasField();
// exit(EXIT_SUCCESS);
// Set initial slice-to-vol parameters. This option is only for testing/debugging/personal use
if (clo.IsSliceToVol() && clo.InitS2VFname() != std::string("")) {
sm.SetMovementModelOrder(clo.MovementModelOrder(0));
if (clo.RegisterDWI() && clo.Registerb0()) sm.SetS2VMovement(clo.InitS2VFname(),ScanType::Any);
else if (clo.RegisterDWI()) sm.SetS2VMovement(clo.InitS2VFname(),ScanType::DWI);
else sm.SetS2VMovement(clo.InitS2VFname(),ScanType::b0);
}
//////////////////////////////////////////////////////////////////////
// Next do the slice-to-vol registration if requested.
//////////////////////////////////////////////////////////////////////
if (clo.IsSliceToVol() && clo.S2V_NIter(0) > 0 && clo.InitS2VFname() == std::string("")) {
double s2v_key = prof.StartEntry("Calling DoSliceToVolumeRegistration");
if (clo.Verbose()) cout << "Performing slice-to-volume registration" << endl;
if (clo.EstimateMoveBySusc()) { // Restrict EC if we are to eventually estimate MBS
SecondLevelECModelType b0_slm = clo.b0_SecondLevelModel();
SecondLevelECModelType dwi_slm = clo.SecondLevelModel();
clo.Set_b0_SecondLevelModel(SecondLevelECModelType::Linear);
clo.SetSecondLevelModel(SecondLevelECModelType::Linear);
for (unsigned int i=0; i<clo.NumOfNonZeroMovementModelOrder(); i++) {
if (clo.Verbose()) cout << "Setting slice-to-volume order to " << clo.MovementModelOrder(i) << endl;
sm.SetMovementModelOrder(clo.MovementModelOrder(i));
sm.Set_S2V_Lambda(clo.S2V_Lambda(i));
dwi_rm = DoSliceToVolumeRegistration(clo,i,false,sm,dwi_rm);
sm.ApplyLocationReference();
}
clo.Set_b0_SecondLevelModel(b0_slm);
clo.SetSecondLevelModel(dwi_slm);
}
else {
for (unsigned int i=0; i<clo.NumOfNonZeroMovementModelOrder(); i++) {
if (clo.Verbose()) cout << "Setting slice-to-volume order to " << clo.MovementModelOrder(i) << endl;
sm.SetMovementModelOrder(clo.MovementModelOrder(i));
sm.Set_S2V_Lambda(clo.S2V_Lambda(i));
dwi_rm = DoSliceToVolumeRegistration(clo,i,false,sm,dwi_rm);
sm.ApplyLocationReference();
}
}
prof.EndEntry(s2v_key);
// Do another check for residual position differences in case s2v wrecked things
if (clo.RegisterDWI()) {
double peas_key = prof.StartEntry("Calling PostEddyAlignShells after S2V");
if (!original_sep_offs_move && !clo.AlignShellsPostEddy()) {
if (clo.Verbose()) cout << "Checking shell alignment along PE-direction (running PostEddyAlignShellsAlongPE)" << endl;
PEASUtils::PostEddyAlignShellsAlongPE(clo,false,sm);
if (clo.Verbose()) cout << "Checking shell alignment (running PostEddyAlignShells)" << endl;
PEASUtils::PostEddyAlignShells(clo,false,sm);
}
else if (original_sep_offs_move && !clo.AlignShellsPostEddy()) {
if (clo.Verbose()) cout << "Aligning shells along PE-direction (running PostEddyAlignShellsAlongPE)" << endl;
PEASUtils::PostEddyAlignShellsAlongPE(clo,true,sm);
if (clo.Verbose()) cout << "Checking shell alignment (running PostEddyAlignShells)" << endl;
PEASUtils::PostEddyAlignShells(clo,false,sm);
}
else if (clo.AlignShellsPostEddy()) {
if (clo.Verbose()) cout << "Checking shell alignment along PE-direction (running PostEddyAlignShellsAlongPE)" << endl;
PEASUtils::PostEddyAlignShellsAlongPE(clo,false,sm);
if (clo.Verbose()) cout << "Aligning shells (running PostEddyAlignShells)" << endl;
PEASUtils::PostEddyAlignShells(clo,true,sm);
}
prof.EndEntry(peas_key);
}
}
// Set initial MBS derivative fields. This option is only for testing/debugging/personal use
if (clo.EstimateMoveBySusc() && clo.InitMBSFname() != std::string("")) {
NEWIMAGE::volume4D<float> mbs_init;
NEWIMAGE::read_volume4D(mbs_init,clo.InitMBSFname());
for (unsigned int i=0; i<clo.MoveBySuscParam().size(); i++) {
sm.SetDerivSuscField(clo.MoveBySuscParam()[i],mbs_init[i]);
}
}
//////////////////////////////////////////////////////////////////////
// Estimate movement-by-susceptibility interaction if requested.
//////////////////////////////////////////////////////////////////////
if (clo.EstimateMoveBySusc() && clo.MoveBySuscNiter()>0) {
sm.SetUseB0sToInformDWIRegistration(false); // So as to preserve current estimates
std::vector<unsigned int> b0s;
std::vector<unsigned int> dwis;
if (clo.IsSliceToVol()) { // Use only "steady" volumes if we know who they are
EDDY::s2vQuant s2vq(sm,1.0,1.0); // "Steadiness hardcoded to 1mm and 1degree
if (clo.Registerb0()) b0s = s2vq.FindStillVolumes(ScanType::b0,clo.MoveBySuscParam());
if (clo.RegisterDWI()) dwis = s2vq.FindStillVolumes(ScanType::DWI,clo.MoveBySuscParam());
}
else { // Otherwise use all volumes
if (clo.Registerb0()) { b0s.resize(sm.NScans(ScanType::b0)); for (unsigned int i=0; i<sm.NScans(ScanType::b0); i++) b0s[i] = i; }
if (clo.RegisterDWI()) { dwis.resize(sm.NScans(ScanType::DWI)); for (unsigned int i=0; i<sm.NScans(ScanType::DWI); i++) dwis[i] = i; }
}
if (clo.RegisterDWI()) { // Do interleaved MBS and EC/movement estimation
unsigned int mbs_niter = (clo.MoveBySuscNiter() / clo.N_MBS_Interleaves()) + 1;
unsigned int niter, s2vi=0;
if (clo.IsSliceToVol()) {
s2vi = clo.NumOfNonZeroMovementModelOrder()-1;
niter = clo.S2V_NIter(s2vi);
sm.SetMovementModelOrder(clo.MovementModelOrder(s2vi));
sm.Set_S2V_Lambda(clo.S2V_Lambda(s2vi));
clo.SetS2VParam(clo.MovementModelOrder(s2vi),clo.S2V_Lambda(s2vi),0.0,(niter/clo.N_MBS_Interleaves())+1);
}
else { niter = clo.NIter(); clo.SetNIterAndFWHM((niter/clo.N_MBS_Interleaves())+1,std::vector<float>(1,0.0)); }
NEWMAT::ColumnVector spar;
EDDY::MoveBySuscCF cf(sm,clo,b0s,dwis,clo.MoveBySuscParam(),clo.MoveBySuscOrder(),clo.MoveBySuscKsp());
for (unsigned int i=0; i<clo.N_MBS_Interleaves(); i++) {
if (clo.Verbose()) cout << "Running interleave " << i+1 << " of MBS" << endl;
if (!i) spar = cf.Par(); // Start guesses all zeros for first iteration
cf.SetLambda(clo.MoveBySuscLambda());
MISCMATHS::NonlinParam nlp(cf.NPar(),MISCMATHS::NL_LM,spar);
nlp.SetMaxIter(mbs_niter);
nlp.SetGaussNewtonType(MISCMATHS::LM_GN);
double mbs_key = prof.StartEntry("Calling nonlin for MBS");
MISCMATHS::nonlin(nlp,cf);
prof.EndEntry(mbs_key);
spar = cf.Par(); // Save for next iteration
if (clo.IsSliceToVol()) {
if (clo.Verbose()) cout << "Running slice-to-vol interleaved with MBS" << endl;
double s2v_key = prof.StartEntry("Calling DoSliceToVolumeRegistration as part of MBS");
dwi_rm = DoSliceToVolumeRegistration(clo,s2vi,false,sm,dwi_rm);
sm.ApplyLocationReference();
prof.EndEntry(s2v_key);
}
else {
if (clo.Verbose()) cout << "Running vol-to-vol interleaved with MBS" << endl;
double v2v_key = prof.StartEntry("Calling DoVolumeToVolumeRegistration as part of MBS");
dwi_rm = DoVolumeToVolumeRegistration(clo,sm);
sm.ApplyLocationReference();
prof.EndEntry(v2v_key);
}
}
cf.WriteFirstOrderFields(clo.MoveBySuscFirstOrderFname());
if (clo.MoveBySuscOrder() > 1) cf.WriteSecondOrderFields(clo.MoveBySuscSecondOrderFname());
}
else { // Just do a straightforward MBS estimation
// Make cost-function object for movement-by-susceptibility
EDDY::MoveBySuscCF cf(sm,clo,b0s,dwis,clo.MoveBySuscParam(),clo.MoveBySuscOrder(),clo.MoveBySuscKsp());
NEWMAT::ColumnVector spar = cf.Par(); // Start guesses (all zeros);
cf.SetLambda(clo.MoveBySuscLambda());
MISCMATHS::NonlinParam nlp(cf.NPar(),MISCMATHS::NL_LM,spar);
nlp.SetMaxIter(clo.MoveBySuscNiter());
nlp.SetGaussNewtonType(MISCMATHS::LM_GN);
double mbs_key = prof.StartEntry("Calling nonlin for MBS");
MISCMATHS::nonlin(nlp,cf); // Estimate
prof.EndEntry(mbs_key);
cf.WriteFirstOrderFields(clo.MoveBySuscFirstOrderFname());
if (clo.MoveBySuscOrder() > 1) cf.WriteSecondOrderFields(clo.MoveBySuscSecondOrderFname());
}
// Do another check for residual position differences in case MBS wrecked things
if (clo.RegisterDWI()) {
double peas_key = prof.StartEntry("Calling PostEddyAlignShells after MBS");
if (!original_sep_offs_move && !clo.AlignShellsPostEddy()) {
if (clo.Verbose()) cout << "Checking shell alignment along PE-direction (running PostEddyAlignShellsAlongPE)" << endl;
PEASUtils::PostEddyAlignShellsAlongPE(clo,false,sm);
if (clo.Verbose()) cout << "Checking shell alignment (running PostEddyAlignShells)" << endl;
PEASUtils::PostEddyAlignShells(clo,false,sm);
}
else if (original_sep_offs_move && !clo.AlignShellsPostEddy()) {
if (clo.Verbose()) cout << "Aligning shells along PE-direction (running PostEddyAlignShellsAlongPE)" << endl;
PEASUtils::PostEddyAlignShellsAlongPE(clo,true,sm);
if (clo.Verbose()) cout << "Checking shell alignment (running PostEddyAlignShells)" << endl;
PEASUtils::PostEddyAlignShells(clo,false,sm);
}
else if (clo.AlignShellsPostEddy()) {
if (clo.Verbose()) cout << "Checking shell alignment along PE-direction (running PostEddyAlignShellsAlongPE)" << endl;
PEASUtils::PostEddyAlignShellsAlongPE(clo,false,sm);
if (clo.Verbose()) cout << "Aligning shells (running PostEddyAlignShells)" << endl;
PEASUtils::PostEddyAlignShells(clo,true,sm);
}
prof.EndEntry(peas_key);
}
}
//////////////////////////////////////////////////////////////////////
// Do final check (and possibly replacement) of outliers with ff=1
// (very little Q-space smoothing).
//////////////////////////////////////////////////////////////////////
if (clo.RegisterDWI()) {
if (clo.Verbose()) cout << "Performing final outlier check" << endl;
double old_hypar_ff = 1.0;
if (clo.HyParFudgeFactor() != 1.0) { old_hypar_ff = clo.HyParFudgeFactor(); clo.SetHyParFudgeFactor(1.0); }
double folc_key = prof.StartEntry("Calling FinalOLCheck");
dwi_rm = FinalOLCheck(clo,dwi_rm,sm);
prof.EndEntry(folc_key);
if (old_hypar_ff != 1.0) clo.SetHyParFudgeFactor(old_hypar_ff);
// Write outlier information
double wol_key = prof.StartEntry("Writing outlier information");
std::vector<unsigned int> i2i = sm.GetDwi2GlobalIndexMapping();
json_out.outlier_report = dwi_rm->WriteReport(i2i,clo.OLReportFname(),!clo.SupressOldStyleTextFiles());
json_out.outlier_maps = dwi_rm->WriteMatrixReport(i2i,sm.NScans(),clo.OLMapReportFname(),clo.OLNStDevMapReportFname(),
clo.OLNSqrStDevMapReportFname(),!clo.SupressOldStyleTextFiles());
if (clo.WriteOutlierFreeData()) {
if (clo.Verbose()) cout << "Running sm.WriteOutlierFreeData" << endl;
sm.WriteOutlierFreeData(clo.OLFreeDataFname());
}
prof.EndEntry(wol_key);
}
// Add rotation. Hidden function. ONLY to be used to
// test that rotation of b-vecs does the right thing.
if (clo.DoTestRot()) {
if (clo.Verbose()) cout << "Running sm.AddRotation" << endl;
sm.AddRotation(clo.TestRotAngles());
}
// Write registration parameters
if (clo.Verbose()) cout << "Running sm.WriteParameterFile" << endl;
if (clo.RegisterDWI() && clo.Registerb0())
json_out.parameters = sm.WriteParameterFile(clo.ParOutFname(),ScanType::Any,!clo.SupressOldStyleTextFiles());
else if (clo.RegisterDWI())
json_out.parameters = sm.WriteParameterFile(clo.ParOutFname(),ScanType::DWI,!clo.SupressOldStyleTextFiles());
else json_out.parameters = sm.WriteParameterFile(clo.ParOutFname(),ScanType::b0,!clo.SupressOldStyleTextFiles());
if (clo.EstimateLongEC()) {
if (clo.Verbose()) cout << "Running sm.LongTimeConstantECModel().WriteReport" << endl;
json_out.long_ec_parameters = sm.LongTimeConstantECModel().WriteReport(sm,clo.LongECOutFname(),!clo.SupressOldStyleTextFiles());
}
// Write movement-over-time file if SliceToVol registration was performed
if (sm.IsSliceToVol()) {
if (clo.Verbose()) cout << "Running sm.WriteMovementOverTimeFile" << endl;
if (clo.RegisterDWI() && clo.Registerb0())
json_out.movement_over_time = sm.WriteMovementOverTimeFile(clo.MovementOverTimeOutFname(),ScanType::Any,!clo.SupressOldStyleTextFiles());
else if (clo.RegisterDWI())
json_out.movement_over_time = sm.WriteMovementOverTimeFile(clo.MovementOverTimeOutFname(),ScanType::DWI,!clo.SupressOldStyleTextFiles());
else json_out.movement_over_time = sm.WriteMovementOverTimeFile(clo.MovementOverTimeOutFname(),ScanType::b0,!clo.SupressOldStyleTextFiles());
}
// Write registered images
if (clo.Verbose()) cout << "Running sm.WriteRegisteredImages" << endl;
double wri_key = prof.StartEntry("Writing registered images");
if (!clo.ReplaceOutliers()) { if (clo.Verbose()) { cout << "Running sm.RecycleOutliers" << endl; } sm.RecycleOutliers(); } // Bring back original data
if (sm.IsSliceToVol()) { // If we need to get predictions to support the resampling
NEWIMAGE::volume4D<float> pred;
ScanType st;
if (clo.RegisterDWI() && clo.Registerb0()) st=ScanType::Any; else if (clo.RegisterDWI()) st=ScanType::DWI; else st=ScanType::b0;
GetPredictionsForResampling(clo,st,sm,pred);
// Set resampling in slice direction to spline.
EDDY::PolationPara old_pp = sm.GetPolation();
EDDY::PolationPara new_pp = old_pp;
if (old_pp.GetS2VInterp() != NEWIMAGE::spline) { new_pp.SetS2VInterp(NEWIMAGE::spline); sm.SetPolation(new_pp); }
if (clo.RegisterDWI() && clo.Registerb0()) sm.WriteRegisteredImages(clo.IOutFname(),clo.OutMaskFname(),clo.ResamplingMethod(),clo.LSResamplingLambda(),clo.MaskOutput(),pred,clo.NumberOfThreads());
else if (clo.RegisterDWI()) sm.WriteRegisteredImages(clo.IOutFname(),clo.OutMaskFname(),clo.ResamplingMethod(),clo.LSResamplingLambda(),clo.MaskOutput(),pred,clo.NumberOfThreads(),ScanType::DWI);
else sm.WriteRegisteredImages(clo.IOutFname(),clo.OutMaskFname(),clo.ResamplingMethod(),clo.LSResamplingLambda(),clo.MaskOutput(),pred,clo.NumberOfThreads(),ScanType::b0);
if (old_pp.GetS2VInterp() != NEWIMAGE::spline) sm.SetPolation(old_pp);
}
else {
if (clo.RegisterDWI() && clo.Registerb0()) sm.WriteRegisteredImages(clo.IOutFname(),clo.OutMaskFname(),clo.ResamplingMethod(),clo.LSResamplingLambda(),clo.MaskOutput(),clo.NumberOfThreads());
else if (clo.RegisterDWI()) sm.WriteRegisteredImages(clo.IOutFname(),clo.OutMaskFname(),clo.ResamplingMethod(),clo.LSResamplingLambda(),clo.MaskOutput(),clo.NumberOfThreads(),ScanType::DWI);
else sm.WriteRegisteredImages(clo.IOutFname(),clo.OutMaskFname(),clo.ResamplingMethod(),clo.LSResamplingLambda(),clo.MaskOutput(),clo.NumberOfThreads(),ScanType::b0);
}
prof.EndEntry(wri_key);
// Optionally write out a "data set" that consists of the GP predictions. This
// was added to demonstrate to Kings that it makes b***er all difference.
if (clo.WritePredictions() || clo.WriteScatterBrainPredictions()) {
// First write predictions in model space
NEWIMAGE::volume4D<float> pred;
EDDY::ScanType st;
std::vector<double> hypar;
if (clo.RegisterDWI() && clo.Registerb0()) st=ScanType::Any; else if (clo.RegisterDWI()) st=ScanType::DWI; else st=ScanType::b0;
if (clo.WritePredictions()) { // Do the "regular" fitting to the GP, but ensure spline interpolation in z if s2v
if (clo.Verbose()) cout << "Running EDDY::GetPredictionsForResampling" << endl;
double wpred_key = prof.StartEntry("Calculating and writing predictions");
EDDY::PolationPara old_pol = sm.GetPolation();
EDDY::PolationPara new_pol = old_pol;
if (clo.IsSliceToVol() && new_pol.GetS2VInterp() != NEWIMAGE::spline) new_pol.SetS2VInterp(NEWIMAGE::spline);
sm.SetPolation(new_pol);
EddyCommandLineOptions tmp_clo = clo;
if (!tmp_clo.RotateBVecsDuringEstimation()) tmp_clo.SetRotateBVecsDuringEstimation(true);
hypar = GetPredictionsForResampling(tmp_clo,st,sm,pred);
pred /= sm.ScaleFactor();
NEWIMAGE::write_volume(pred,clo.PredictionsOutFname());
sm.SetPolation(old_pol);
// Next write predictions in scan space(s)
pred = 0.0;
hypar = GetPredictionsForResampling(clo,st,sm,pred);
pred /= sm.ScaleFactor();
for (int s=0; s<pred.tsize(); s++) {
pred[s] = EddyUtils::TransformModelToScanSpace(pred[s],sm.Scan(s,st),sm.GetSuscHzOffResField(s,st));
}
NEWIMAGE::write_volume(pred,clo.PredictionsInScanSpaceOutFname());
prof.EndEntry(wpred_key);
}
if (clo.WriteScatterBrainPredictions()) { // Write predictions using "scattered data approach"
if (clo.Verbose()) cout << "Running EDDY::GetScatterBrainPredictions" << endl;
double wspred_key = prof.StartEntry("Calculating and writing scatter brain predictions");
GetScatterBrainPredictions(clo,st,sm,hypar,pred);
pred /= sm.ScaleFactor();
NEWIMAGE::write_volume(pred,clo.ScatterBrainPredictionsOutFname());
prof.EndEntry(wspred_key);
}
if (clo.WriteVolumeScatterBrainPredictions()) { // Write predictions using "scattered data approach" but with volume based bvec rotation
if (clo.Verbose()) cout << "Running EDDY::GetScatterBrainPredictions with volume based bvec rotation" << endl;
double wspred_key = prof.StartEntry("Calculating and writing scatter brain predictions");
GetScatterBrainPredictions(clo,st,sm,hypar,pred,true);
pred /= sm.ScaleFactor();
NEWIMAGE::write_volume(pred,clo.VolumeScatterBrainPredictionsOutFname());
prof.EndEntry(wspred_key);
}
}
// Optionally write an additional set of registered images, this time with outliers retained.
// This was added for the benefit of the HCP.
if (clo.ReplaceOutliers() && clo.WriteAdditionalResultsWithOutliersRetained()) {
double wri_key = prof.StartEntry("Writing registered images with outliers retained");
if (clo.Verbose()) cout << "Running sm.WriteRegisteredImages" << endl;
if (clo.Verbose()) { cout << "Running sm.RecycleOutliers" << endl; }
sm.RecycleOutliers(); // Bring back original data
if (sm.IsSliceToVol()) { // If we need to get predictions to support the resampling
NEWIMAGE::volume4D<float> pred;
ScanType st;
if (clo.RegisterDWI() && clo.Registerb0()) st=ScanType::Any; else if (clo.RegisterDWI()) st=ScanType::DWI; else st=ScanType::b0;
GetPredictionsForResampling(clo,st,sm,pred);
if (clo.RegisterDWI() && clo.Registerb0()) sm.WriteRegisteredImages(clo.AdditionalWithOutliersOutFname(),std::string(""),clo.ResamplingMethod(),clo.LSResamplingLambda(),clo.MaskOutput(),pred,clo.NumberOfThreads());
else if (clo.RegisterDWI()) sm.WriteRegisteredImages(clo.AdditionalWithOutliersOutFname(),std::string(""),clo.ResamplingMethod(),clo.LSResamplingLambda(),clo.MaskOutput(),pred,clo.NumberOfThreads(),ScanType::DWI);
else sm.WriteRegisteredImages(clo.AdditionalWithOutliersOutFname(),std::string(""),clo.ResamplingMethod(),clo.LSResamplingLambda(),clo.MaskOutput(),pred,clo.NumberOfThreads(),ScanType::b0);
}
else {
if (clo.RegisterDWI() && clo.Registerb0()) sm.WriteRegisteredImages(clo.AdditionalWithOutliersOutFname(),std::string(""),clo.ResamplingMethod(),clo.LSResamplingLambda(),clo.MaskOutput(),clo.NumberOfThreads());
else if (clo.RegisterDWI()) sm.WriteRegisteredImages(clo.AdditionalWithOutliersOutFname(),std::string(""),clo.ResamplingMethod(),clo.LSResamplingLambda(),clo.MaskOutput(),clo.NumberOfThreads(),ScanType::DWI);
else sm.WriteRegisteredImages(clo.AdditionalWithOutliersOutFname(),std::string(""),clo.ResamplingMethod(),clo.LSResamplingLambda(),clo.MaskOutput(),clo.NumberOfThreads(),ScanType::b0);
}
prof.EndEntry(wri_key);
}
// Write EC fields
if (clo.WriteFields()) {
if (clo.Verbose()) cout << "Running sm.WriteECFields" << endl;
if (clo.RegisterDWI() && clo.Registerb0()) sm.WriteECFields(clo.ECFOutFname());
else if (clo.RegisterDWI()) sm.WriteECFields(clo.ECFOutFname(),ScanType::DWI);
else sm.WriteECFields(clo.ECFOutFname(),ScanType::b0);
}
// Write rotated b-vecs
if (clo.WriteRotatedBVecs()) {
if (clo.Verbose()) cout << "Running sm.WriteRotatedBVecs" << endl;
sm.WriteRotatedBVecs(clo.RotatedBVecsOutFname());
}
// Write b-vecs for LSR resampling
if (clo.ResamplingMethod() == FinalResamplingType::LSR) {
if (clo.Verbose()) cout << "Running sm.WriteBVecsForLSR" << endl;
sm.WriteBVecsForLSR(clo.BVecsLSROutFname(),false);
if (clo.WriteRotatedBVecs()) {
if (clo.Verbose()) cout << "Running sm.WriteBVecsForLSR with rotation" << endl;
sm.WriteBVecsForLSR(clo.RotatedBVecsLSROutFname(),true);
}
}
// Write movement RMS
if (clo.WriteMovementRMS()) {
double rms_key = prof.StartEntry("Writing RMS");
if (clo.Verbose()) cout << "Running sm.WriteMovementRMS" << endl;
if (clo.RegisterDWI() && clo.Registerb0()) {
json_out.movement_rms = sm.WriteMovementRMS(clo.RMSOutFname(),ScanType::Any,!clo.SupressOldStyleTextFiles());
json_out.restricted_movement_rms = sm.WriteRestrictedMovementRMS(clo.RestrictedRMSOutFname(),ScanType::Any,!clo.SupressOldStyleTextFiles());
}
else if (clo.RegisterDWI()) {
json_out.movement_rms = sm.WriteMovementRMS(clo.RMSOutFname(),ScanType::DWI,!clo.SupressOldStyleTextFiles());
json_out.restricted_movement_rms = sm.WriteRestrictedMovementRMS(clo.RestrictedRMSOutFname(),ScanType::DWI,!clo.SupressOldStyleTextFiles());
}
else {
json_out.movement_rms = sm.WriteMovementRMS(clo.RMSOutFname(),ScanType::b0,!clo.SupressOldStyleTextFiles());
json_out.restricted_movement_rms = sm.WriteRestrictedMovementRMS(clo.RestrictedRMSOutFname(),ScanType::b0,!clo.SupressOldStyleTextFiles());
}
prof.EndEntry(rms_key);
}
// Write the .json file
json_out.Write(clo.JsonOutFname());
// Write CNR maps
if (clo.WriteCNRMaps() || clo.WriteRangeCNRMaps() || clo.WriteResiduals()) {
double cnr_key = prof.StartEntry("Writing CNR maps");
double old_hypar_ff = 1.0;
if (clo.HyParFudgeFactor() != 1.0) { old_hypar_ff = clo.HyParFudgeFactor(); clo.SetHyParFudgeFactor(1.0); }
if (clo.Verbose()) cout << "Running EDDY::WriteCNRMaps" << endl;
WriteCNRMaps(clo,sm,clo.CNROutFname(),clo.RangeCNROutFname(),clo.ResidualsOutFname());
if (old_hypar_ff != 1.0) clo.SetHyParFudgeFactor(old_hypar_ff);
prof.EndEntry(cnr_key);
}
// Write 3D displacement fields
if (clo.WriteDisplacementFields()) {
if (clo.Verbose()) cout << "Running sm.WriteDisplacementFields" << endl;
if (clo.RegisterDWI() && clo.Registerb0()) sm.WriteDisplacementFields(clo.DFieldOutFname());
else if (clo.RegisterDWI()) sm.WriteDisplacementFields(clo.DFieldOutFname(),ScanType::DWI);
else sm.WriteDisplacementFields(clo.DFieldOutFname(),ScanType::b0);
}
prof.EndEntry(total_key);
return(EXIT_SUCCESS);
}
catch(const std::exception& e)
{
cout << "EDDY::: Eddy failed with message " << e.what() << endl;
return(EXIT_FAILURE);
}
catch(...)
{
cout << "EDDY::: Eddy failed" << endl;
return(EXIT_FAILURE);
}
namespace EDDY {
/****************************************************************//**
*
* A very high-level global function that registers all the scans
* (b0 and dwis) in sm using a volume-to-volume model.
* \param[in] clo Carries information about the command line options
* that eddy was invoked with.
* \param[in,out] sm Collection of all scans. Will be updated by this call.
* \return A ptr to ReplacementManager that details which slices in
* which dwi scans were replaced by their expectations.
*
********************************************************************/
ReplacementManager *DoVolumeToVolumeRegistration(// Input
const EddyCommandLineOptions& clo,
// Input/Output
ECScanManager& sm) EddyTry
{
static Utilities::FSLProfiler prof("_"+string(__FILE__)+"_"+string(__func__));
double total_key = prof.StartEntry("Total");
if (clo.IsDiffusion()) {
// Start by registering the b0 scans
NEWMAT::Matrix b0_mss, b0_ph;
ReplacementManager *b0_rm = nullptr;
if (clo.NIter() && clo.Registerb0() && sm.NScans(ScanType::b0)>1) {
double b0_key = prof.StartEntry("b0");
if (clo.Verbose()) cout << "Running Register" << endl;
b0_rm = Register(clo,ScanType::b0,clo.NIter(),clo.FWHM(),clo.b0_SecondLevelModel(),false,sm,b0_rm,b0_mss,b0_ph);
if (clo.IsSliceToVol()) { // Set scan for shape reference if we are to do also slice-to-volume registration
if (clo.ShapeReferencesSet()) { // User has specified shape-references
if (clo.Verbose()) cout << "Setting user requested scan " << clo.B0ShapeReference() << " as b0 shape-reference."<< endl;
sm.SetB0ShapeReference(clo.B0ShapeReference());
}
else { // We need to find best scan for shape reference
double minmss=1e20;
unsigned int mindx=0;
for (unsigned int i=0; i<sm.NScans(ScanType::b0); i++) {
if (b0_mss(b0_mss.Nrows(),i+1) < minmss) { minmss=b0_mss(b0_mss.Nrows(),i+1); mindx=i; }
}
if (clo.Verbose()) cout << "Setting scan " << sm.Getb02GlobalIndexMapping(mindx) << " as b0 shape-reference."<< endl;
sm.SetB0ShapeReference(sm.Getb02GlobalIndexMapping(mindx));
}
}
// Apply reference for location
if (clo.Verbose()) cout << "Running sm.ApplyB0LocationReference" << endl;
sm.ApplyB0LocationReference();
prof.EndEntry(b0_key);
}
// See if we can use b0 movement estimates to inform dwi registration
if (sm.B0sAreInterspersed() && sm.UseB0sToInformDWIRegistration() && clo.Registerb0() && clo.RegisterDWI() && clo.NIter() > 0) {
if (clo.Verbose()) cout << "Running sm.PolateB0MovPar" << endl;
sm.PolateB0MovPar();
}
// Now register the dwi scans
NEWMAT::Matrix dwi_mss, dwi_ph;
ReplacementManager *dwi_rm = NULL;
if (clo.NIter() && clo.RegisterDWI()) {
double dwi_key = prof.StartEntry("dwi");
if (clo.Verbose()) cout << "Running Register" << endl;
dwi_rm = Register(clo,ScanType::DWI,clo.NIter(),clo.FWHM(),clo.SecondLevelModel(),true,sm,dwi_rm,dwi_mss,dwi_ph);
if (clo.IsSliceToVol()) { // Set scan for shape reference if we are to do also slice-to-volume registration
if (clo.ShapeReferencesSet()) { // User has specified shape-reference
for (unsigned int i=0; i<sm.NoOfShells(ScanType::DWI); i++) {
std::pair<int,double> ref = clo.ShellShapeReference(i);
if (clo.Verbose()) cout << "Setting user requested scan " << ref.first << " as shell shape-reference for shell "<< i << " with b-value= " << ref.second << endl;
sm.SetShellShapeReference(i,ref.first);
}
}
else { // We need to find best scan for shape reference
std::vector<double> bvals;
std::vector<std::vector<unsigned int> > shindx = sm.GetShellIndicies(bvals);
for (unsigned int shell=0; shell<shindx.size(); shell++) {
double minmss=1e20;
unsigned int mindx=0;
bool found_vol_with_no_outliers=false;
for (unsigned int i=0; i<shindx[shell].size(); i++) {
if (!sm.Scan(shindx[shell][i]).HasOutliers()) { // Only consider scans without outliers
found_vol_with_no_outliers=true;
if (dwi_mss(dwi_mss.Nrows(),sm.GetGlobal2DWIIndexMapping(shindx[shell][i])+1) < minmss) {
minmss=dwi_mss(dwi_mss.Nrows(),sm.GetGlobal2DWIIndexMapping(shindx[shell][i])+1);
mindx=shindx[shell][i];
}
}
}
if (!found_vol_with_no_outliers) {
std::vector<unsigned int> i2i = sm.GetDwi2GlobalIndexMapping();
dwi_rm->WriteReport(i2i,clo.OLReportFname());
dwi_rm->WriteMatrixReport(i2i,sm.NScans(),clo.OLMapReportFname(true),clo.OLNStDevMapReportFname(true),clo.OLNSqrStDevMapReportFname(true));
std::ostringstream errtxt;
errtxt << "DoVolumeToVolumeRegistration: Unable to find volume with no outliers in shell " << shell << " with b-value=" << bvals[shell];
throw EddyException(errtxt.str());
}
if (clo.Verbose()) cout << "Setting scan " << mindx << " as shell shape-reference for shell "<< shell << " with b-value= " << bvals[shell] << endl;
sm.SetShellShapeReference(shell,mindx);
}
}
// Apply reference for location
}
if (clo.Verbose()) cout << "Running sm.ApplyDWILocationReference" << endl;
sm.ApplyDWILocationReference();
prof.EndEntry(dwi_key);
}
// Write history of cost-function and parameter estimates if requested
if (clo.NIter() && clo.History()) {
if (clo.RegisterDWI()) {
MISCMATHS::write_ascii_matrix(clo.DwiMssHistoryFname(),dwi_mss);
MISCMATHS::write_ascii_matrix(clo.DwiParHistoryFname(),dwi_ph);
}
if (clo.Registerb0()) {
MISCMATHS::write_ascii_matrix(clo.B0MssHistoryFname(),b0_mss);
MISCMATHS::write_ascii_matrix(clo.B0ParHistoryFname(),b0_ph);
}
}
prof.EndEntry(total_key);
return(dwi_rm);
}
else { // If it is fMRI
// Now register the dwi scans
NEWMAT::Matrix fmri_mss, fmri_ph;
ReplacementManager *fmri_rm = NULL;
if (clo.NIter()) {
double fmri_key = prof.StartEntry("fmri");
if (clo.Verbose()) cout << "Running Register" << endl;
fmri_rm = Register(clo,ScanType::fMRI,clo.NIter(),clo.FWHM(),clo.SecondLevelModel(),true,sm,fmri_rm,fmri_mss,fmri_ph);
if (clo.IsSliceToVol()) { // Set scan for shape reference if we are to do also slice-to-volume registration
if (clo.ShapeReferencesSet()) { // User has specified shape-references
if (clo.Verbose()) cout << "Setting user requested scan " << clo.fMRIShapeReference() << " as shape-reference."<< endl;
sm.SetfMRIShapeReference(clo.fMRIShapeReference());
}
else { // We need to find best scan for shape reference
double minmss=1e20;
unsigned int mindx=0;
bool found_vol_with_no_outliers=false;
for (unsigned int i=0; i<sm.NScans(ScanType::fMRI); i++) {
if (!sm.Scan(i,ScanType::fMRI).HasOutliers()) {
found_vol_with_no_outliers=true;
if (fmri_mss(fmri_mss.Nrows(),i+1) < minmss) {
minmss = fmri_mss(fmri_mss.Nrows(),i+1);
mindx = i;
}
}
}
if (!found_vol_with_no_outliers) {
std::vector<unsigned int> one2one(sm.NScans(ScanType::fMRI));
std::iota(one2one.begin(),one2one.end(),0);
fmri_rm->WriteReport(one2one,clo.OLReportFname());
fmri_rm->WriteMatrixReport(one2one,sm.NScans(),clo.OLMapReportFname(true),clo.OLNStDevMapReportFname(true),clo.OLNSqrStDevMapReportFname(true));
throw EddyException("DoVolumeToVolumeRegistration: Unable to find volume with no outliers.");
}
if (clo.Verbose()) cout << "Setting scan " << mindx << " as shape-reference."<< endl;
sm.SetfMRIShapeReference(mindx);
}
// Apply reference for location
}
if (clo.Verbose()) cout << "Running sm.ApplyDWILocationReference" << endl;
sm.ApplyfMRILocationReference();
prof.EndEntry(fmri_key);
}
// Write history of cost-function and parameter estimates if requested
if (clo.NIter() && clo.History()) {
MISCMATHS::write_ascii_matrix(clo.fMRIMssHistoryFname(),fmri_mss);
MISCMATHS::write_ascii_matrix(clo.fMRIParHistoryFname(),fmri_ph);
}
prof.EndEntry(total_key);
return(fmri_rm);
}
} EddyCatch
/****************************************************************//**
*
* A very high-level global function that estimates long
* time-constant eddy currents. It will alternate between estimating
* the long EC parameters and updating the EC parameters.
* \param[in] clo Carries information about the command line options
* that eddy was invoked with.
* \param[in,out] sm Collection of all scans. Will be updated by this call.
* \param[in,out] rm A ptr to ReplacementManager that details which slices
* in which dwi scans were replaced by their expectations.
* \return A ptr to ReplacementManager that details which slices in
* which dwi scans were replaced by their expectations.
*
********************************************************************/
ReplacementManager *EstimateLongEC(// Input
const EddyCommandLineOptions& clo,
// Input/Output
ECScanManager& sm,
ReplacementManager *rm) EddyTry
{
static Utilities::FSLProfiler prof("_"+string(__FILE__)+"_"+string(__func__));
double total_key = prof.StartEntry("Total");
if (sm.IsDiffusion()) {
NEWIMAGE::volume<float> mask = sm.Scan(0).GetIma(); EddyUtils::SetTrilinearInterp(mask); mask = 1.0; // FOV mask in model space
for (unsigned int iter=0; iter<clo.LongECNIter(); iter++) {
#ifdef COMPILE_GPU
// Load prediction maker in model space and make all the predictions
// Note that for this we need to have predictions both for b0 and dwi
std::shared_ptr<DWIPredictionMaker> b0_pmp = EddyGpuUtils::LoadPredictionMaker(clo,ScanType::b0,sm,iter,0.0,mask);
std::shared_ptr<DWIPredictionMaker> dwi_pmp = EddyGpuUtils::LoadPredictionMaker(clo,ScanType::DWI,sm,iter,0.0,mask);
std::vector<NEWIMAGE::volume<float> > pred(sm.NScans());
for (GpuPredictorChunk c(sm,ScanType::b0); c<sm.NScans(ScanType::b0); c++) {
std::vector<unsigned int> b0i = c.Indicies();
std::vector<NEWIMAGE::volume<float> > chunk = b0_pmp->Predict(b0i);
for (const unsigned int& i : b0i) pred[sm.Getb02GlobalIndexMapping(i)] = chunk[i];
}
for (GpuPredictorChunk c(sm,ScanType::DWI); c<sm.NScans(ScanType::DWI); c++) {
std::vector<unsigned int> dwii = c.Indicies();
std::vector<NEWIMAGE::volume<float> > chunk = dwi_pmp->Predict(dwii);
for (const unsigned int& i : dwii) pred[sm.GetDwi2GlobalIndexMapping(i)] = chunk[i];
}
// Now estimate long EC parameters
if (clo.DebugLevel()) {
std::vector<double> mss = EddyGpuUtils::LongECParamUpdate(pred,mask,0.0,clo.VeryVerbose(),iter,clo.DebugLevel(),clo.DebugIndicies().GetIndicies(),sm);
}
else {
std::vector<double> mss = EddyGpuUtils::LongECParamUpdate(pred,mask,0.0,clo.VeryVerbose(),sm);
}
#else
// Load prediction maker in model space and make all the predictions
// Note that for this we need to have predictions both for b0 and dwi
std::shared_ptr<DWIPredictionMaker> b0_pmp = EDDY::LoadPredictionMaker(clo,ScanType::b0,sm,iter,0.0,mask);
std::shared_ptr<DWIPredictionMaker> dwi_pmp = EDDY::LoadPredictionMaker(clo,ScanType::DWI,sm,iter,0.0,mask);
std::vector<NEWIMAGE::volume<float> > pred(sm.NScans());
for (unsigned int s=0; s<sm.NScans(ScanType::b0); s++) pred[sm.Getb02GlobalIndexMapping(s)] = b0_pmp->Predict(s);
for (unsigned int s=0; s<sm.NScans(ScanType::DWI); s++) pred[sm.GetDwi2GlobalIndexMapping(s)] = dwi_pmp->Predict(s);
// Caclulate parameter updates
if (clo.DebugLevel()) {
std::vector<double> mss = EddyUtils::LongECParamUpdate(pred,mask,0.0,clo.NumberOfThreads(),clo.VeryVerbose(),iter,clo.DebugLevel(),clo.DebugIndicies().GetIndicies(),sm);
}
else {
std::vector<double> mss = EddyUtils::LongECParamUpdate(pred,mask,0.0,clo.NumberOfThreads(),clo.VeryVerbose(),sm);
}
#endif
if (clo.ReestimateECWhenEstimatingLongEC()) {
// Next update the "regular" movement and/or EC parameters
ReplacementManager *b0_rm = nullptr;
NEWMAT::Matrix b0_mss, b0_ph;
b0_rm = Register(clo,ScanType::b0,1,std::vector<float>(1,0.0),clo.b0_SecondLevelModel(),false,sm,b0_rm,b0_mss,b0_ph);
NEWMAT::Matrix dwi_mss, dwi_ph;
rm = Register(clo,ScanType::DWI,1,std::vector<float>(1,0.0),clo.SecondLevelModel(),true,sm,rm,dwi_mss,dwi_ph);
}
}
}
else throw EddyException("EstimateLongEC: Called with fMRI data");
prof.EndEntry(total_key);
return(rm);
} EddyCatch
/****************************************************************//**
*
* A very high-level global function that registers all the scans
* (b0 and dwis) in sm using a slice-to-volume model.
* \param[in] clo Carries information about the command line options
* that eddy was invoked with.
* \param[in,out] sm Collection of all scans. Will be updated by this call.
* \return A ptr to ReplacementManager that details which slices in
* which dwi scans were replaced by their expectations.
*
********************************************************************/
ReplacementManager *DoSliceToVolumeRegistration(// Input
const EddyCommandLineOptions& clo,
unsigned int oi, // Order index
bool dol, // Detect outliers?
// Input/Output
ECScanManager& sm,
ReplacementManager *dwi_rm) EddyTry
{
static Utilities::FSLProfiler prof("_"+string(__FILE__)+"_"+string(__func__));
double total_key = prof.StartEntry("Total");
if (sm.IsDiffusion()) {
// Start by registering the b0 scans
NEWMAT::Matrix b0_mss, b0_ph;
ReplacementManager *b0_rm = NULL;
if (clo.S2V_NIter(oi) && clo.Registerb0() && sm.NScans(ScanType::b0)>1) {
double b0_key = prof.StartEntry("b0");
if (clo.Verbose()) cout << "Running Register" << endl;
b0_rm = Register(clo,ScanType::b0,clo.S2V_NIter(oi),clo.S2V_FWHM(oi),clo.b0_SecondLevelModel(),false,sm,b0_rm,b0_mss,b0_ph);
// Set reference for shape
if (clo.Verbose()) cout << "Running sm.ApplyB0ShapeReference" << endl;
sm.ApplyB0ShapeReference();
// Set reference for location
if (clo.Verbose()) cout << "Running sm.ApplyB0LocationReference" << endl;
sm.ApplyB0LocationReference();
prof.EndEntry(b0_key);
}
// Now register the dwi scans
NEWMAT::Matrix dwi_mss, dwi_ph;
if (clo.S2V_NIter(oi) && clo.RegisterDWI()) {
double dwi_key = prof.StartEntry("dwi");
if (clo.Verbose()) cout << "Running Register" << endl;
dwi_rm = Register(clo,ScanType::DWI,clo.S2V_NIter(oi),clo.S2V_FWHM(oi),clo.SecondLevelModel(),dol,sm,dwi_rm,dwi_mss,dwi_ph);
// Set reference for shape
if (clo.Verbose()) cout << "Running sm.ApplyShellShapeReference" << endl;
for (unsigned int si=0; si<sm.NoOfShells(ScanType::DWI); si++) sm.ApplyShellShapeReference(si);
// Set reference for location
if (clo.Verbose()) cout << "Running sm.ApplyDWILocationReference" << endl;
sm.ApplyDWILocationReference();
prof.EndEntry(dwi_key);
}
// Write history of cost-function and parameter estimates if requested
if (clo.S2V_NIter(oi) && clo.History()) {
if (clo.RegisterDWI()) {
MISCMATHS::write_ascii_matrix(clo.DwiMssS2VHistoryFname(),dwi_mss);
MISCMATHS::write_ascii_matrix(clo.DwiParS2VHistoryFname(),dwi_ph);
}
if (clo.Registerb0()) {
MISCMATHS::write_ascii_matrix(clo.B0MssS2VHistoryFname(),b0_mss);
MISCMATHS::write_ascii_matrix(clo.B0ParS2VHistoryFname(),b0_ph);
}
}
prof.EndEntry(total_key);
return(dwi_rm);
}
else { // Assume fMRI
NEWMAT::Matrix fmri_mss, fmri_ph;
ReplacementManager *fmri_rm = NULL;
double fmri_key = prof.StartEntry("fmri");
if (clo.Verbose()) cout << "Running Register" << endl;
fmri_rm = Register(clo,ScanType::fMRI,clo.S2V_NIter(oi),clo.S2V_FWHM(oi),clo.b0_SecondLevelModel(),false,sm,fmri_rm,fmri_mss,fmri_ph);
// Set reference for shape
if (clo.Verbose()) cout << "Running sm.ApplyfMRIShapeReference" << endl;
sm.ApplyfMRIShapeReference();
// Set reference for location
if (clo.Verbose()) cout << "Running sm.ApplyfMRILocationReference" << endl;
sm.ApplyfMRILocationReference();
prof.EndEntry(fmri_key);
// Write history of cost-function and parameter estimates if requested
if (clo.S2V_NIter(oi) && clo.History()) {
MISCMATHS::write_ascii_matrix(clo.fMRIMssS2VHistoryFname(),fmri_mss);
MISCMATHS::write_ascii_matrix(clo.fMRIParS2VHistoryFname(),fmri_ph);
}
return(fmri_rm);
}
} EddyCatch
/****************************************************************//**
*
* A very high-level global function that estimates a receieve
* bias field that is stationary in the scanner framework.
* \param[in] clo Carries information about the command line options
* that eddy was invoked with.
* \param[in] ksp Knotspacing of spline representation of field.
* Set to negative value for free-form field.
* \param[in] lambda Weight of regularisation when estimating field.
* \param[in,out] sm Collection of all scans. Will be updated by this call.
*
********************************************************************/
/*
void EstimateBiasField(// Input
const EddyCommandLineOptions& clo,
double ksp,
double lambda,
// Input/output
ECScanManager& sm) EddyTry
{
BiasFieldEstimator bfe;
ScanType st[2] = {ScanType::b0, ScanType::DWI};
NEWIMAGE::volume<float> mask = sm.Scan(0,ScanType::Any).GetIma(); EddyUtils::SetTrilinearInterp(mask); mask = 1.0; // FOV-mask in model space
#ifdef COMPILE_GPU
// Loop over b0s and DWIs
for (unsigned int sti=0; sti<2; sti++) {
if (sm.NScans(st[sti]) > 1) { // Will only have info on field if 2 or more locations
std::shared_ptr<DWIPredictionMaker> pmp = EddyGpuUtils::LoadPredictionMaker(clo,st[sti],sm,0,0.0,mask);
// First calculate the average location for this scan type to use as reference
EDDY::ImageCoordinates mic = sm.Scan(0,st[sti]).SamplingPoints();
for (unsigned int s=1; s<sm.NScans(st[sti]); s++) {
mic += sm.Scan(s,st[sti]).SamplingPoints();
}
mic /= static_cast<float>(sm.NScans(st[sti]));
bfe.SetRefScan(mask,mic);
// Next loop over chunks of predictions and add scans
for (GpuPredictorChunk c(sm,st[sti]); c<sm.NScans(st[sti]); c++) {
std::vector<unsigned int> si = c.Indicies();
std::vector<NEWIMAGE::volume<float> > pred = pmp->Predict(si);
// Loop over scans within chunk
for (unsigned int i=0; i<si.size(); i++) {
// EDDY::ImageCoordinates ic = EddyGpuUtils::GetCoordinatesInScannerSpace(sm.Scan(si[i],st[sti]));
// cout << "i = " << i << ", calling SamplingPoints()" << endl;
EDDY::ImageCoordinates ic = sm.Scan(si[i],st[sti]).SamplingPoints();
bfe.AddScan(pred[i],EddyGpuUtils::GetUnwarpedScan(sm.Scan(si[i],st[sti]),sm.GetSuscHzOffResField(si[i],st[sti]),true,sm.GetPolation()),mask,ic);
}
}
}
}
#else
// Loop over b0s and DWIs
for (unsigned int sti=0; sti<2; sti++) {
std::shared_ptr<DWIPredictionMaker> pmp = EDDY::LoadPredictionMaker(clo,st[sti],sm,0,0.0,mask);
// Loop over scans
#pragma omp parallel for
for (int s=0; s<int(sm.NScans(st[sti])); s++) {
// Get coordinates in scanner space
EDDY::ImageCoordinates ic = sm.Scan(s,st[sti]).SamplingPoints();
// Get prediction in model space
NEWIMAGE::volume<float> pred = pmp->Predict(s);
// Get original data in model space
NEWIMAGE::volume<float> sm.Scan(s,st[sti]).GetUnwarpedIma(sm.GetSuscHzOffResField(s,st[sti]));
// AddScan
bfe.AddScan(pred,sm.Scan(s,st[sti]).GetUnwarpedIma(sm.GetSuscHzOffResField(s,st[sti])),mask,ic);
}
}
#endif
// Caclulate many fields
// double iksp[] = {10.0};
// double ilambda[] = {1e-10, 1e-6, 1e-2, 100, 1e6, 1e10, 1e14};
// for (unsigned int i=0; i<7; i++) {
// cout << "Calculating field with lambda = " << ilambda[i] << endl;
// NEWIMAGE::volume<float> bfield = bfe.GetField(ilambda[i]);
// char fname[256]; sprintf(fname,"bfield_%d",i);
// NEWIMAGE::write_volume(bfield,fname);
// }
// Calculate field
NEWIMAGE::volume<float> bfield = bfe.GetField(1e10);
NEWIMAGE::write_volume(bfield,"bfield_1e10");
// Set field
sm.SetBiasField(bfield);
// Do a second step where the unknown offset is calculated by
// maximising the SNR/CNR of the corrected data.
// float offset = EstimateBiasFieldOffset();
return;
} EddyCatch
*/
/****************************************************************//**
*
* A very high-level global function that estimates the offset
* of an estimated bias field. It does that by finding the offset
* that maximises a weighted sum of the SNR (b0-volumes) and CNR
* (dwi-volumes) of the corrected data. It will return the estimated
* offset, and also set it in the ScanManager. It will also set the
* average of the bias field to one within the user defined mask.
* \returns offset The estimated offset value
* \param[in] clo Carries information about the command line options
* that eddy was invoked with.
* \param[in,out] sm Collection of all scans. Will be updated by this call.
*
********************************************************************/
float EstimatedBiasFieldOffset(// Input
const EddyCommandLineOptions& clo,
// Input/output
ECScanManager& sm) EddyTry
{
return(1.0);
} EddyCatch
/****************************************************************//**
*
* A global function that registers the scans in sm.
* \param[in] clo Carries information about the command line options
* that eddy was invoked with.
* \param[in] st Specifies if we should register the diffusion weighted
* images or the b=0 images.
* \param[in] niter Specifies how many iterations should be run
* \param[in] slm Specifies if a 2nd level model should be used to
* constrain the estimates.
* \param[in] dol Detect outliers if true
* \param[in,out] sm Collection of all scans. Will be updated by this call.
* \param[in,out] rm Pointer to ReplacementManager. If NULL on input
* one will be allocated and the new pointer value passed on return.
* \param[out] msshist Returns the history of the mss. msshist(i,j)
* contains the mss for the jth scan on the ith iteration.
* \param[out] phist Returns the history of the estimated parameters.
* phist(i,j) contains the jth parameter on the ith iteration
* \return A ptr to ReplacementManager that details which slices in
* which scans were replaced by their expectations.
*
********************************************************************/
ReplacementManager *Register(// Input
const EddyCommandLineOptions& clo,
ScanType st,
unsigned int niter,
const std::vector<float>& fwhm,
SecondLevelECModelType slm,
bool dol,
// Input/Output
ECScanManager& sm,
ReplacementManager *rm,
// Output
NEWMAT::Matrix& msshist,
NEWMAT::Matrix& phist) EddyTry
{
static Utilities::FSLProfiler prof("_"+string(__FILE__)+"_"+string(__func__));
double total_key = prof.StartEntry("Total");
msshist.ReSize(niter,sm.NScans(st));
phist.ReSize(niter,sm.NScans(st)*sm.Scan(0,st).NParam());
double *mss_tmp = new double[sm.NScans(st)]; // Replace by vector
if (rm == NULL) { // If no replacement manager passed in
if (clo.OLSummaryStatistics() == OLSumStats::ShellWise) {
std::vector<double> bvals;
std::vector<std::vector<unsigned int> > shi = sm.GetDWIShellIndicies(bvals);
rm = new ReplacementManager(shi,bvals,static_cast<unsigned int>(sm.Scan(0,st).GetIma().zsize()),clo.OLDef(),clo.OLErrorType(),clo.OLType(),clo.MultiBand());
}
else rm = new ReplacementManager(sm.NScans(st),static_cast<unsigned int>(sm.Scan(0,st).GetIma().zsize()),clo.OLDef(),clo.OLErrorType(),clo.OLType(),clo.MultiBand());
}
NEWIMAGE::volume<float> mask = sm.Scan(0,st).GetIma(); EddyUtils::SetTrilinearInterp(mask); mask = 1.0; // FOV-mask in model space
for (unsigned int iter=0; iter<niter; iter++) {
double load_key = prof.StartEntry("LoadPredictionMaker");
// Load prediction maker in model space
#ifdef COMPILE_GPU
std::shared_ptr<DWIPredictionMaker> pmp = EddyGpuUtils::LoadPredictionMaker(clo,st,sm,iter,fwhm[iter],mask);
if (clo.DebugLevel() > 2) EDDY::WritePMDebugInfo(pmp,clo,st,iter,"GPU_","before_repol");
#else
std::shared_ptr<DWIPredictionMaker> pmp = EDDY::LoadPredictionMaker(clo,st,sm,iter,fwhm[iter],mask);
if (clo.DebugLevel() > 2) EDDY::WritePMDebugInfo(pmp,clo,st,iter,"","before_repol");
#endif
prof.EndEntry(load_key);
// Detect outliers and replace them
DiffStatsVector stats(sm.NScans(st));
if (dol) {
double dol_key = prof.StartEntry("Detecting and replacing outliers");
std::shared_ptr<DWIPredictionMaker> od_pmp;
#ifdef COMPILE_GPU
od_pmp = pmp;
if (clo.DebugLevel()) stats = EddyGpuUtils::DetectOutliers(clo,st,od_pmp,mask,sm,iter,clo.DebugLevel(),*rm);
else stats = EddyGpuUtils::DetectOutliers(clo,st,od_pmp,mask,sm,*rm);
if (iter) {
EddyGpuUtils::ReplaceOutliers(clo,st,od_pmp,mask,*rm,false,sm);
// EddyGpuUtils::UpdatePredictionMaker(clo,st,sm,rm,mask,pmp);
pmp = EddyGpuUtils::LoadPredictionMaker(clo,st,sm,iter,fwhm[iter],mask);
if (clo.DebugLevel() > 2) EDDY::WritePMDebugInfo(pmp,clo,st,iter,"GPU_","after_repol");
}
#else
od_pmp = pmp;
if (clo.DebugLevel()) stats = EDDY::DetectOutliers(clo,st,od_pmp,mask,sm,iter,clo.DebugLevel(),*rm);
else stats = EDDY::DetectOutliers(clo,st,od_pmp,mask,sm,*rm);
if (iter) {
EDDY::ReplaceOutliers(clo,st,od_pmp,mask,*rm,false,sm);
// EDDY::UpdatePredictionMaker(clo,st,sm,rm,mask,pmp);
pmp = EDDY::LoadPredictionMaker(clo,st,sm,iter,fwhm[iter],mask);
if (clo.DebugLevel() > 2) EDDY::WritePMDebugInfo(pmp,clo,st,iter,"","after_repol");
}
#endif
prof.EndEntry(dol_key);
}
//
// Calculate the parameter updates
// Note that this section will proceed in three different
// ways depending on if it is run in single processor mode,
// multi processor mode (OpenMP) or with a GPU-box (CUDA).
//
if (clo.Verbose()) cout << "Calculating parameter updates" << endl;
double update_key = prof.StartEntry("Updating parameters");
#ifdef COMPILE_GPU
for (GpuPredictorChunk c(sm,st); c<sm.NScans(st); c++) {
double predict_chunk_key = prof.StartEntry("Predicting chunk");
std::vector<unsigned int> si = c.Indicies();
if (clo.VeryVerbose()) cout << "Making predictions for scans: " << c << endl;
std::vector<NEWIMAGE::volume<float> > pred = pmp->Predict(si);
prof.EndEntry(predict_chunk_key);
if (clo.VeryVerbose()) cout << "Finished making predictions for scans: " << c << endl;
double update_chunk_key = prof.StartEntry("Updating parameters for chunk");
for (unsigned int i=0; i<si.size(); i++) {
unsigned int global_indx = sm.GetGlobalIndex(si[i],st);
if (clo.DebugLevel() && clo.DebugIndicies().IsAmongIndicies(global_indx)) {
mss_tmp[si[i]] = EddyGpuUtils::MovAndECParamUpdate(pred[i],sm.GetSuscHzOffResField(si[i],st),sm.GetBiasField(),mask,fwhm[iter],clo.VeryVerbose(),global_indx,iter,clo.DebugLevel(),sm.Scan(si[i],st));
}
else mss_tmp[si[i]] = EddyGpuUtils::MovAndECParamUpdate(pred[i],sm.GetSuscHzOffResField(si[i],st),sm.GetBiasField(),mask,fwhm[iter],clo.VeryVerbose(),global_indx,sm.Scan(si[i],st));
if (clo.VeryVerbose()) {
std::cout << "Iter: " << iter << ", scan: " << global_indx << ", gpu_mss = " << mss_tmp[si[i]] << std::endl << std::flush;
}
}
prof.EndEntry(update_chunk_key);
}
#else
if (clo.NumberOfThreads() == 1) { // If we are to run single threaded
for (int s=0; s<int(sm.NScans(st)); s++) {
// Get prediction in model space
NEWIMAGE::volume<float> pred = pmp->Predict(s);
// Update parameters
unsigned int global_indx = sm.GetGlobalIndex(s,st);
if (clo.DebugLevel() && clo.DebugIndicies().IsAmongIndicies(global_indx)) {
mss_tmp[s] = EddyUtils::MovAndECParamUpdate(pred,sm.GetSuscHzOffResField(s,st),sm.GetBiasField(),mask,fwhm[iter],clo.VeryVerbose(),global_indx,iter,clo.DebugLevel(),sm.Scan(s,st));
}
else mss_tmp[s] = EddyUtils::MovAndECParamUpdate(pred,sm.GetSuscHzOffResField(s,st),sm.GetBiasField(),mask,fwhm[iter],clo.VeryVerbose(),sm.Scan(s,st));
if (clo.VeryVerbose()) {
std::cout << "Iter: " << iter << ", scan: " << global_indx << ", mss = " << mss_tmp[s] << std::endl << std::flush;
}
}
}
else { // We are to run multi-threaded
// Decide number of volumes per thread
std::vector<unsigned int> nvols = EddyUtils::ScansPerThread(sm.NScans(st),clo.NumberOfThreads());
// Next spawn threads to do the calculations
std::vector<std::thread> threads(clo.NumberOfThreads()-1); // + main thread makes clo.NumberOfThreads()
for (unsigned int i=0; i<clo.NumberOfThreads()-1; i++) {
threads[i] = std::thread(MovAndECParamUpdateWrapper,nvols[i],nvols[i+1],std::ref(clo),
pmp,std::ref(mask),st,fwhm[iter],iter,std::ref(sm),mss_tmp);
}
MovAndECParamUpdateWrapper(nvols[clo.NumberOfThreads()-1],nvols[clo.NumberOfThreads()],clo,pmp,mask,st,fwhm[iter],iter,sm,mss_tmp);
std::for_each(threads.begin(),threads.end(),std::mem_fn(&std::thread::join));
}
#endif
prof.EndEntry(update_key);
// Print/collect some information that can be used for diagnostics
Diagnostics(clo,iter,st,sm,mss_tmp,stats,*rm,msshist,phist);
// Maybe use model based EC parameters
if (slm != SecondLevelECModelType::None) {
if (clo.VeryVerbose()) cout << "Performing 2nd level modelling of estimated parameters" << endl;
sm.SetPredictedECParam(st,slm);
}
// Maybe try and separate field-offset and translation in PE direction
if (clo.SeparateOffsetFromMovement()) {
if (clo.VeryVerbose()) cout << "Attempting to separate field-offset from subject movement" << endl;
sm.SeparateFieldOffsetFromMovement(st,clo.OffsetModel());
}
}
delete [] mss_tmp;
prof.EndEntry(total_key);
return(rm);
} EddyCatch
ReplacementManager *FinalOLCheck(// Input
const EddyCommandLineOptions& clo,
// Input/output
ReplacementManager *rm,
ECScanManager& sm) EddyTry
{
static Utilities::FSLProfiler prof("_"+string(__FILE__)+"_"+string(__func__));
double total_key = prof.StartEntry("Total");
NEWIMAGE::volume<float> mask = sm.Scan(0,ScanType::DWI).GetIma(); EddyUtils::SetTrilinearInterp(mask); mask = 1.0; // FOV-mask in model space
if (rm == NULL) {
if (clo.OLSummaryStatistics() == OLSumStats::ShellWise) {
std::vector<double> bvals;
std::vector<std::vector<unsigned int> > shi = sm.GetDWIShellIndicies(bvals);
rm = new ReplacementManager(shi,bvals,static_cast<unsigned int>(sm.Scan(0,ScanType::DWI).GetIma().zsize()),clo.OLDef(),clo.OLErrorType(),clo.OLType(),clo.MultiBand());
}
rm = new ReplacementManager(sm.NScans(ScanType::DWI),static_cast<unsigned int>(sm.Scan(0,ScanType::DWI).GetIma().zsize()),clo.OLDef(),clo.OLErrorType(),clo.OLType(),clo.MultiBand());
}
// Load prediction maker in model space
double load_key = prof.StartEntry("LoadPredictionMaker");
#ifdef COMPILE_GPU
std::shared_ptr<DWIPredictionMaker> pmp = EddyGpuUtils::LoadPredictionMaker(clo,ScanType::DWI,sm,0,0.0,mask);
#else
std::shared_ptr<DWIPredictionMaker> pmp = EDDY::LoadPredictionMaker(clo,ScanType::DWI,sm,0,0.0,mask);
#endif
prof.EndEntry(load_key);
// Detect outliers and replace them
DiffStatsVector stats(sm.NScans(ScanType::DWI));
bool add_noise = clo.AddNoiseToReplacements();
double det_key = prof.StartEntry("DetectOutliers");
#ifdef COMPILE_GPU
stats = EddyGpuUtils::DetectOutliers(clo,ScanType::DWI,pmp,mask,sm,*rm);
#else
stats = EDDY::DetectOutliers(clo,ScanType::DWI,pmp,mask,sm,*rm);
#endif
prof.EndEntry(det_key);
double rep_key = prof.StartEntry("ReplaceOutliers");
#ifdef COMPILE_GPU
EddyGpuUtils::ReplaceOutliers(clo,ScanType::DWI,pmp,mask,*rm,add_noise,sm);
#else
EDDY::ReplaceOutliers(clo,ScanType::DWI,pmp,mask,*rm,add_noise,sm);
#endif
prof.EndEntry(rep_key);
prof.EndEntry(total_key);
return(rm);
} EddyCatch
/****************************************************************//**
*
* A global function that loads up a prediction maker with all scans
* of a given type. It will load it with unwarped scans (given the
* current estimates of the warps) as served up by sm.GetUnwarpedScan()
* or sm.GetUnwarpedOrigScan() depending on the value of use_orig.
* \param[in] clo Carries information about the command line options
* that eddy was invoked with.
* \param[in] st Specifies if we should register the diffusion weighted
* images or the b=0 images. If it is set to DWI the function will return
* an EDDY::DiffusionGP prediction maker and if it is set to b0 it will
* return an EDDY::b0Predictor.
* \param[in] sm Collection of all scans.
* \param[out] mask Returns a mask that indicates the voxels where data
* is present for all input scans in sm.
* \param[in] use_orig If set to true it will load it with unwarped "original"
* , i.e. un-smoothed, scans. Default is false.
* \return A safe pointer to a DWIPredictionMaker that can be used to
* make predictions about what the scans should look like in undistorted space.
*
********************************************************************/
std::shared_ptr<DWIPredictionMaker> LoadPredictionMaker(// Input
const EddyCommandLineOptions& clo,
ScanType st,
const ECScanManager& sm,
unsigned int iter,
float fwhm,
// Output
NEWIMAGE::volume<float>& mask,
// Optional input
bool use_orig) EddyTry
{
static Utilities::FSLProfiler prof("_"+string(__FILE__)+"_"+string(__func__));
double total_key = prof.StartEntry("Total");
std::shared_ptr<DWIPredictionMaker> pmp; // Prediction Maker Pointer
std::shared_ptr<HyParCF> hpcf; // Hyper-parameter cost-function
std::shared_ptr<HyParEstimator> hpe; // Hyper-parameter estimator
if (st==ScanType::DWI || (st==ScanType::fMRI && !clo.HyperParFixed())) { // If we actually need to estimate hyper-parameters
if (clo.HyperParFixed()) hpe = std::shared_ptr<FixedValueHyParEstimator>(new FixedValueHyParEstimator(clo.HyperParValues()));
else {
if (clo.HyParCostFunction() == HyParCostFunctionType::CC) hpe = std::shared_ptr<CheapAndCheerfulHyParEstimator>(new CheapAndCheerfulHyParEstimator(clo.NVoxHp(),clo.InitRand()));
else {
if (clo.HyParCostFunction() == HyParCostFunctionType::MML) hpcf = std::shared_ptr<MMLHyParCF>(new MMLHyParCF);
else if (clo.HyParCostFunction() == HyParCostFunctionType::CV) hpcf = std::shared_ptr<CVHyParCF>(new CVHyParCF);
else if (clo.HyParCostFunction() == HyParCostFunctionType::GPP) hpcf = std::shared_ptr<GPPHyParCF>(new GPPHyParCF);
else throw EddyException("LoadPredictionMaker: Unknown hyperparameter cost-function");
hpe = std::shared_ptr<FullMontyHyParEstimator>(new FullMontyHyParEstimator(hpcf,clo.HyParFudgeFactor(),clo.NVoxHp(),clo.InitRand(),clo.VeryVerbose()));
}
}
}
if (st==ScanType::DWI) { // If diffusion weighted data
std::shared_ptr<KMatrix> K;
if (clo.CovarianceFunction() == CovarianceFunctionType::Spherical) K = std::shared_ptr<SphericalKMatrix>(new SphericalKMatrix(clo.DontCheckShelling()));
else if (clo.CovarianceFunction() == CovarianceFunctionType::Exponential) K = std::shared_ptr<ExponentialKMatrix>(new ExponentialKMatrix(clo.DontCheckShelling()));
else if (clo.CovarianceFunction() == CovarianceFunctionType::NewSpherical) K = std::shared_ptr<NewSphericalKMatrix>(new NewSphericalKMatrix(clo.DontCheckShelling()));
else throw EddyException("LoadPredictionMaker: Unknown covariance function");
pmp = std::shared_ptr<DWIPredictionMaker>(new DiffusionGP(K,hpe)); // GP
}
else if (st==ScanType::b0) pmp = std::shared_ptr<DWIPredictionMaker>(new b0Predictor); // Silly mean predictor for b=0 data
else if (st==ScanType::fMRI) {
if (clo.CovarianceFunction() == CovarianceFunctionType::SquaredExponential) {
if (clo.HyperParFixed()) pmp = std::shared_ptr<DWIPredictionMaker>(new b0Predictor); // Use mean predictor for now
else {
std::shared_ptr<KMatrix> K = std::shared_ptr<SqrExpKMatrix>(new SqrExpKMatrix());
pmp = std::shared_ptr<DWIPredictionMaker>(new fmriPredictor(K,hpe));
}
}
else throw EddyException("LoadPredictionMaker: Unknown covariance function");
}
else throw EddyException("LoadPredictionMaker: Invalid scan type");
pmp->SetNoOfScans(sm.NScans(st));
mask = sm.Scan(0,st).GetIma(); EddyUtils::SetTrilinearInterp(mask); mask = 1.0;
double load_key = prof.StartEntry("Loading");
if (clo.Verbose()) std::cout << "Loading prediction maker" << std::flush;
if (clo.VeryVerbose()) std::cout << std::endl << "Scan:" << std::flush;
if (clo.NumberOfThreads() == 1) { // If we are to run single threaded
for (int s=0; s<int(sm.NScans(st)); s++) {
if (clo.VeryVerbose()) std::cout << " " << sm.GetGlobalIndex(s,st) << std::flush;
NEWIMAGE::volume<float> tmpmask = sm.Scan(s,st).GetIma();
EddyUtils::SetTrilinearInterp(tmpmask); tmpmask = 1.0;
if (use_orig) pmp->SetScan(sm.GetUnwarpedOrigScan(s,tmpmask,st),sm.Scan(s,st).GetDiffPara(clo.RotateBVecsDuringEstimation()),s);
else pmp->SetScan(sm.GetUnwarpedScan(s,tmpmask,st),sm.Scan(s,st).GetDiffPara(clo.RotateBVecsDuringEstimation()),s);
{
mask *= tmpmask;
}
}
if (clo.VeryVerbose()) std::cout << std::endl << std::flush;
}
else { // We are to run multi-threaded
// Decide number of volumes per thread
std::vector<unsigned int> nvols = EddyUtils::ScansPerThread(sm.NScans(st),clo.NumberOfThreads());
// Next spawn threads to do the calculations
std::vector<std::thread> threads(clo.NumberOfThreads()-1); // + main thread makes clo.NumberOfThreads()
for (unsigned int i=0; i<clo.NumberOfThreads()-1; i++) {
threads[i] = std::thread(SetUnwarpedScanWrapper,nvols[i],nvols[i+1],std::ref(clo),
std::ref(sm),st,use_orig,pmp,std::ref(mask));
}
SetUnwarpedScanWrapper(nvols[clo.NumberOfThreads()-1],nvols[clo.NumberOfThreads()],clo,sm,st,use_orig,pmp,mask);
std::for_each(threads.begin(),threads.end(),std::mem_fn(&std::thread::join));
if (clo.VeryVerbose()) std::cout << std::endl << std::flush;
}
prof.EndEntry(load_key);
double eval_key = prof.StartEntry("Evaluating");
if (clo.Verbose()) cout << endl << "Evaluating prediction maker model" << endl;
pmp->EvaluateModel(sm.Mask()*mask,fwhm,clo.Verbose());
prof.EndEntry(eval_key);
if (clo.DebugLevel() > 2 && st==ScanType::DWI) {
char fname[256];
sprintf(fname,"EDDY_DEBUG_K_Mat_Data_%02d",iter);
pmp->WriteMetaData(fname);
}
prof.EndEntry(total_key);
return(pmp);
} EddyCatch
DiffStatsVector DetectOutliers(// Input
const EddyCommandLineOptions& clo,
ScanType st,
std::shared_ptr<const DWIPredictionMaker> pmp,
const NEWIMAGE::volume<float>& mask,
const ECScanManager& sm,
// Input/Output
ReplacementManager& rm) EddyTry
{
return(detect_outliers(clo,st,pmp,mask,sm,0,0,rm));
} EddyCatch
DiffStatsVector DetectOutliers(// Input
const EddyCommandLineOptions& clo,
ScanType st,
std::shared_ptr<const DWIPredictionMaker> pmp,
const NEWIMAGE::volume<float>& mask,
const ECScanManager& sm,
unsigned int iter,
unsigned int dl, // Debug Level
// Input/Output
ReplacementManager& rm) EddyTry
{
return(detect_outliers(clo,st,pmp,mask,sm,iter,dl,rm));
} EddyCatch
DiffStatsVector detect_outliers(// Input
const EddyCommandLineOptions& clo,
ScanType st,
std::shared_ptr<const DWIPredictionMaker> pmp,
const NEWIMAGE::volume<float>& mask,
const ECScanManager& sm,
unsigned int iter,
unsigned int dl,
// Input/Output
ReplacementManager& rm) EddyTry
{
static Utilities::FSLProfiler prof("_"+string(__FILE__)+"_"+string(__func__));
double total_key = prof.StartEntry("Total");
if (clo.VeryVerbose()) cout << "Checking for outliers" << endl;
// Generate slice-wise stats on difference between observation and prediction
DiffStatsVector stats(sm.NScans(st));
if (clo.NumberOfThreads() == 1) { // If we are to run single threaded
for (int s=0; s<int(sm.NScans(st)); s++) {
NEWIMAGE::volume<float> pred = pmp->Predict(s);
stats[s] = EddyUtils::GetSliceWiseStats(clo,st,pred,sm.GetSuscHzOffResField(s,st),mask,sm,s,iter,dl);
}
}
else { // We are running multi-threaded
// Decide number of volumes per thread
std::vector<unsigned int> nvols = EddyUtils::ScansPerThread(sm.NScans(st),clo.NumberOfThreads());
// Next spawn threads to do the calculations
std::vector<std::thread> threads(clo.NumberOfThreads()-1); // + main thread makes clo.NumberOfThreads()
for (unsigned int i=0; i<clo.NumberOfThreads()-1; i++) {
threads[i] = std::thread(GetSliceWiseStatsWrapper,nvols[i],nvols[i+1],std::ref(clo),
std::ref(sm),st,pmp,std::ref(mask),iter,dl,std::ref(stats));
}
GetSliceWiseStatsWrapper(nvols[clo.NumberOfThreads()-1],nvols[clo.NumberOfThreads()],clo,sm,st,pmp,mask,iter,dl,stats);
std::for_each(threads.begin(),threads.end(),std::mem_fn(&std::thread::join));
}
// Detect outliers and update replacement manager
rm.Update(stats);
prof.EndEntry(total_key);
if (dl) { // Note that we use a local "debug level" rather than clo.DebugLevel()
std::string prefix("EDDY_DEBUG_");
if (st==ScanType::b0) prefix += "b0";
else if (st==ScanType::DWI) prefix += "DWI";
else if (st==ScanType::fMRI) prefix += "fMRI";
char fname[256];
sprintf(fname,"%s_repol_info_%02d.mat",prefix.c_str(),iter);
rm.WriteDebugInfo(fname,sm.GetDwi2GlobalIndexMapping(),sm.NScans());
}
return(stats);
} EddyCatch
void ReplaceOutliers(// Input
const EddyCommandLineOptions& clo,
ScanType st,
std::shared_ptr<DWIPredictionMaker> pmp,
const NEWIMAGE::volume<float>& mask,
const ReplacementManager& rm,
bool add_noise,
// Input/Output
ECScanManager& sm) EddyTry
{
static Utilities::FSLProfiler prof("_"+string(__FILE__)+"_"+string(__func__));
double total_key = prof.StartEntry("Total");
// Replace outlier slices with their predictions
if (clo.VeryVerbose()) cout << "Replacing outliers with predictions" << endl;
if (clo.NumberOfThreads() == 1) { // If we are to run single threaded
for (int s=0; s<int(sm.NScans(st)); s++) {
std::vector<unsigned int> ol = rm.OutliersInScan(s);
if (ol.size()) { // If this scan has outlier slices
if (clo.VeryVerbose()) {
std::cout << "Scan " << sm.GetGlobalIndex(s,st) << " has outlier slices:";
for (unsigned int i=0; i<ol.size(); i++) { std::cout << " " << ol[i]; }
std::cout << std::endl << std::flush;
}
NEWIMAGE::volume<float> pred = pmp->Predict(s,true);
if (add_noise) {
double vp = pmp->PredictionVariance(s,true);
double ve = pmp->ErrorVariance(s);
double stdev = std::sqrt(vp+ve) - std::sqrt(vp);
pred += EddyUtils::MakeNoiseIma(pred,0.0,stdev);
}
sm.Scan(s,st).SetAsOutliers(pred,sm.GetSuscHzOffResField(s,st),mask,ol);
}
}
}
else { // We are to run multi-threaded
// Decide number of volumes per thread
std::vector<unsigned int> nvols = EddyUtils::ScansPerThread(sm.NScans(st),clo.NumberOfThreads());
// Next spawn threads to do the calculations
std::vector<std::thread> threads(clo.NumberOfThreads()-1); // + main thread makes clo.NumberOfThreads()
for (unsigned int i=0; i<clo.NumberOfThreads()-1; i++) {
threads[i] = std::thread(SetAsOutliersWrapper,nvols[i],nvols[i+1],pmp,std::ref(mask),
std::ref(rm),st,clo.VeryVerbose(),add_noise,std::ref(sm));
}
SetAsOutliersWrapper(nvols[clo.NumberOfThreads()-1],nvols[clo.NumberOfThreads()],pmp,mask,rm,st,clo.VeryVerbose(),add_noise,sm);
std::for_each(threads.begin(),threads.end(),std::mem_fn(&std::thread::join));
}
prof.EndEntry(total_key);
return;
} EddyCatch
/****************************************************************//**
*
* A global function that makes ff=1 predictions, in model space,
* for all scans of type st.
* \param[in] clo Carries information about the command line options
* that eddy was invoked with.
* \param[in] st Specifies if we should resample the diffusion weighted
* images, the b=0 images or both.
* \param[in] sm Collection of all scans.
* \param[out] pred A 4D volume with predictions for all scans of the
* type indicated by st.
* \return The hyper-parameters used for the predictions
*
********************************************************************/
std::vector<double> GetPredictionsForResampling(// Input
const EddyCommandLineOptions& clo,
ScanType st,
const ECScanManager& sm,
// Output
NEWIMAGE::volume4D<float>& pred) EddyTry
{
static Utilities::FSLProfiler prof("_"+string(__FILE__)+"_"+string(__func__));
double total_key = prof.StartEntry("Total");
pred.reinitialize(sm.Scan(0,st).GetIma().xsize(),sm.Scan(0,st).GetIma().ysize(),sm.Scan(0,st).GetIma().zsize(),sm.NScans(st));
NEWIMAGE::copybasicproperties(sm.Scan(0,st).GetIma(),pred);
NEWIMAGE::volume<float> mask = sm.Scan(0,st).GetIma(); EddyUtils::SetTrilinearInterp(mask); mask = 1.0; // FOV-mask in model space
EddyCommandLineOptions lclo = clo;
std::vector<double> hypar;
if (lclo.HyParFudgeFactor() != 1.0) lclo.SetHyParFudgeFactor(1.0);
if (st == ScanType::Any || st == ScanType::b0) {
#ifdef COMPILE_GPU
std::shared_ptr<DWIPredictionMaker> pmp = EddyGpuUtils::LoadPredictionMaker(lclo,ScanType::b0,sm,0,0.0,mask);
#else
std::shared_ptr<DWIPredictionMaker> pmp = EDDY::LoadPredictionMaker(lclo,ScanType::b0,sm,0,0.0,mask);
#endif
for (unsigned int s=0; s<sm.NScans(ScanType::b0); s++) {
if (st == ScanType::b0) pred[s] = pmp->Predict(s,true);
else pred[sm.Getb02GlobalIndexMapping(s)] = pmp->Predict(s,true);
}
}
if (st == ScanType::Any || st == ScanType::DWI) {
#ifdef COMPILE_GPU
std::shared_ptr<DWIPredictionMaker> pmp = EddyGpuUtils::LoadPredictionMaker(lclo,ScanType::DWI,sm,0,0.0,mask);
#else
std::shared_ptr<DWIPredictionMaker> pmp = EDDY::LoadPredictionMaker(lclo,ScanType::DWI,sm,0,0.0,mask);
#endif
hypar = pmp->GetHyperPar();
for (unsigned int s=0; s<sm.NScans(ScanType::DWI); s++) {
if (st == ScanType::DWI) pred[s] = pmp->Predict(s,true);
else pred[sm.GetDwi2GlobalIndexMapping(s)] = pmp->Predict(s,true);
}
}
prof.EndEntry(total_key);
return(hypar);
} EddyCatch
/****************************************************************//**
*
* A global function that makes ff=1 predictions, in model space,
* for all scans of type st. It differs from the "normal" way eddy
* makes predictions in that it uses a "scattered data reconstruction"
* approach.
* \param[in] clo Carries information about the command line options
* that eddy was invoked with.
* \param[in] st Specifies if we should resample the diffusion weighted
* images, the b=0 images or both.
* \param[in] sm Collection of all scans.
* \param[in] hypar Hyperparameters. Valid only for "spherical" covariance-
* function.
* \param[out] pred A 4D volume with predictions for all scans of the
* type indicated by st.
* \param[in] vwbvrot If true means that bvec rotation is per volume
* instead of per slice/MB-group which is default.
*
********************************************************************/
void GetScatterBrainPredictions(// Input
const EddyCommandLineOptions& clo,
ScanType st,
ECScanManager& sm,
const std::vector<double>& hypar,
// Output
NEWIMAGE::volume4D<float>& pred,
// Optional input
bool vwbvrot) EddyTry
{
static Utilities::FSLProfiler prof("_"+string(__FILE__)+"_"+string(__func__));
double total_key = prof.StartEntry("Total");
// Do some checking to ensure that the call and the input makes sense
if (clo.CovarianceFunction() != CovarianceFunctionType::NewSpherical) {
throw EddyException("EDDY::GetScatterBrainPredictions: Predictions only available for Spherical covariance function");
}
if (!clo.IsSliceToVol()) throw EddyException("EDDY::GetScatterBrainPredictions: Predictions only makes sense for slice-to-vol model");
NEWIMAGE::volume<float> mask = sm.Scan(0,st).GetIma(); EddyUtils::SetTrilinearInterp(mask); mask = 1.0; // FOV-mask in model space
std::vector<DiffPara> dp = sm.GetDiffParas(ScanType::DWI);
NewSphericalKMatrix K(clo.DontCheckShelling());
K.SetDiffusionPar(dp);
if (hypar.size() != K.NoOfHyperPar()) throw EddyException("EDDY::GetScatterBrainPredictions: Incompatible hypar size");
pred.reinitialize(sm.Scan(0,st).GetIma().xsize(),sm.Scan(0,st).GetIma().ysize(),sm.Scan(0,st).GetIma().zsize(),sm.NScans(st));
NEWIMAGE::copybasicproperties(sm.Scan(0,st).GetIma(),pred);
// The b0s predictions are just the means and unaffected by rotations
// Should be updated to be closer to the way we calculate the predictions
// for the diffusion weighted data.
if (st==ScanType::b0 || st==ScanType::Any) {
EDDY::PolationPara old_pol = sm.GetPolation();
EDDY::PolationPara new_pol = old_pol;
if (clo.IsSliceToVol() && new_pol.GetS2VInterp() != NEWIMAGE::trilinear) new_pol.SetS2VInterp(NEWIMAGE::trilinear);
sm.SetPolation(new_pol);
#ifdef COMPILE_GPU
std::shared_ptr<DWIPredictionMaker> pmp = EddyGpuUtils::LoadPredictionMaker(clo,ScanType::b0,sm,0,0.0,mask);
#else
std::shared_ptr<DWIPredictionMaker> pmp = EDDY::LoadPredictionMaker(clo,ScanType::b0,sm,0,0.0,mask);
#endif
for (unsigned int s=0; s<sm.NScans(ScanType::b0); s++) {
if (st == ScanType::b0) pred[s] = pmp->Predict(s,true);
else pred[sm.Getb02GlobalIndexMapping(s)] = pmp->Predict(s,true);
}
sm.SetPolation(new_pol);
}
if (st==ScanType::DWI || st==ScanType::Any) {
#ifdef COMPILE_GPU
EddyGpuUtils::MakeScatterBrainPredictions(clo,sm,hypar,pred,vwbvrot);
#else
throw EddyException("EDDY::GetScatterBrainPredictions: Only implemented for GPU");
#endif
}
prof.EndEntry(total_key);
return;
} EddyCatch
/****************************************************************//**
*
* A global function that calculates CNR-maps for the dwi volume,
* SNR-maps for the b0 volumes and a 4D map of residuals.
* for all scans of type st. All the output parameters are optional
* and passing in a nullptr means that parameter is not calculated
* or returned.
* \param[in] clo Carries information about the command line options
* that eddy was invoked with.
* \param[in] sm Collection of all scans.
* \param[out] std_cnr A 4D file with one CNR-map, calculated as
* std(pred)/std(res), per shell.
* \param[out] range_cnr A 4D file with one CNR-map,calculated as
* range(pred)/std(res), per shell.
* \param[out] b0_snr A 3D SNR-map for the b0-volumes
* \param[out] residuals A 4D-map with the residuals for both dwis and b0s
*
********************************************************************/
void CalculateCNRMaps(// Input
const EddyCommandLineOptions& clo,
const ECScanManager& sm,
// Output
std::shared_ptr<NEWIMAGE::volume4D<float> > std_cnr,
std::shared_ptr<NEWIMAGE::volume4D<float> > range_cnr,
std::shared_ptr<NEWIMAGE::volume<float> > b0_snr,
std::shared_ptr<NEWIMAGE::volume4D<float> > residuals) EddyTry
{
static Utilities::FSLProfiler prof("_"+string(__FILE__)+"_"+string(__func__));
double total_key = prof.StartEntry("Total");
if (std_cnr==nullptr && range_cnr==nullptr && b0_snr==nullptr && residuals==nullptr) {
throw EddyException("EDDY::CalculateCNRMaps: At least one output parameter must be set.");
}
if (!sm.IsShelled()) {
throw EddyException("EDDY::CalculateCNRMaps: Can only calculate CNR for shelled data.");
}
if ((std_cnr!=nullptr && (!NEWIMAGE::samesize(sm.Scan(0,ScanType::Any).GetIma(),(*std_cnr)[0]) || (*std_cnr).tsize() != sm.NoOfShells(ScanType::DWI))) ||
(range_cnr!=nullptr && (!NEWIMAGE::samesize(sm.Scan(0,ScanType::Any).GetIma(),(*range_cnr)[0]) || (*range_cnr).tsize() != sm.NoOfShells(ScanType::DWI))) ||
(b0_snr!=nullptr && !NEWIMAGE::samesize(sm.Scan(0,ScanType::Any).GetIma(),*b0_snr)) ||
(residuals!=nullptr && (!NEWIMAGE::samesize(sm.Scan(0,ScanType::Any).GetIma(),(*residuals)[0]) || (*residuals).tsize() != sm.NScans(ScanType::Any)))) {
throw EddyException("EDDY::CalculateCNRMaps: Size mismatch between sm and output containers.");
}
NEWIMAGE::volume<float> mask = sm.Scan(0,ScanType::DWI).GetIma(); // FOV-mask in model space
EddyUtils::SetTrilinearInterp(mask); mask = 1.0;
std::shared_ptr<DWIPredictionMaker> dwi_pmp;
std::shared_ptr<DWIPredictionMaker> b0_pmp;
if (std_cnr || range_cnr || residuals) {
#ifdef COMPILE_GPU
dwi_pmp = EddyGpuUtils::LoadPredictionMaker(clo,ScanType::DWI,sm,0,0.0,mask);
#else
dwi_pmp = EDDY::LoadPredictionMaker(clo,ScanType::DWI,sm,0,0.0,mask);
#endif
}
if (b0_snr || residuals) {
#ifdef COMPILE_GPU
b0_pmp = EddyGpuUtils::LoadPredictionMaker(clo,ScanType::b0,sm,0,0.0,mask);
#else
b0_pmp = EDDY::LoadPredictionMaker(clo,ScanType::b0,sm,0,0.0,mask);
#endif
}
if (std_cnr || range_cnr) { // Calculate and write shell-wise CNR maps
std::vector<double> grpb; // b-values of the different dwi shells
std::vector<std::vector<unsigned int> > dwi_indx = sm.GetShellIndicies(grpb); // Global indicies of dwi scans
std::vector<NEWIMAGE::volume<float> > mvols(grpb.size()); // Volumes of mean predictions for the shells
// Calculate mean volumes of predictions
for (unsigned int i=0; i<grpb.size(); i++) {
mvols[i] = dwi_pmp->Predict(sm.GetGlobal2DWIIndexMapping(dwi_indx[i][0]));
for (unsigned int j=1; j<dwi_indx[i].size(); j++) {
mvols[i] += dwi_pmp->Predict(sm.GetGlobal2DWIIndexMapping(dwi_indx[i][j]));
}
mvols[i] /= float(dwi_indx[i].size());
}
std::vector<NEWIMAGE::volume<float> > stdvols(grpb.size());
if (std_cnr) { // Calculate standard deviation volumes of predictions
for (unsigned int i=0; i<grpb.size(); i++) {
NEWIMAGE::volume<float> tmp = dwi_pmp->Predict(sm.GetGlobal2DWIIndexMapping(dwi_indx[i][0])) - mvols[i];
stdvols[i] = tmp*tmp;
for (unsigned int j=1; j<dwi_indx[i].size(); j++) {
tmp = dwi_pmp->Predict(sm.GetGlobal2DWIIndexMapping(dwi_indx[i][j])) - mvols[i];
stdvols[i] += tmp*tmp;
}
stdvols[i] /= float(dwi_indx[i].size()-1);
stdvols[i] = NEWIMAGE::sqrt(stdvols[i]);
}
}
std::vector<NEWIMAGE::volume<float> > minvols(grpb.size());
std::vector<NEWIMAGE::volume<float> > maxvols(grpb.size());
if (range_cnr) { // Caclculate range (max-min) volumes of predictions
for (unsigned int i=0; i<grpb.size(); i++) {
minvols[i] = dwi_pmp->Predict(sm.GetGlobal2DWIIndexMapping(dwi_indx[i][0]));
maxvols[i] = minvols[i];
for (unsigned int j=1; j<dwi_indx[i].size(); j++) {
minvols[i] = NEWIMAGE::min(minvols[i],dwi_pmp->Predict(sm.GetGlobal2DWIIndexMapping(dwi_indx[i][j])));
maxvols[i] = NEWIMAGE::max(maxvols[i],dwi_pmp->Predict(sm.GetGlobal2DWIIndexMapping(dwi_indx[i][j])));
}
maxvols[i] -= minvols[i]; // Maxvols now contain the range rather than the max
}
}
// Calculate standard deviation of residuals
std::vector<NEWIMAGE::volume<float> > stdres(grpb.size());
for (unsigned int i=0; i<grpb.size(); i++) {
NEWIMAGE::volume<float> tmp = dwi_pmp->Predict(sm.GetGlobal2DWIIndexMapping(dwi_indx[i][0])) - dwi_pmp->InputData(sm.GetGlobal2DWIIndexMapping(dwi_indx[i][0]));
stdres[i] = tmp*tmp;
for (unsigned int j=1; j<dwi_indx[i].size(); j++) {
tmp = dwi_pmp->Predict(sm.GetGlobal2DWIIndexMapping(dwi_indx[i][j])) - dwi_pmp->InputData(sm.GetGlobal2DWIIndexMapping(dwi_indx[i][j]));
stdres[i] += tmp*tmp;
}
stdres[i] /= float(dwi_indx[i].size()-1);
stdres[i] = NEWIMAGE::sqrt(stdres[i]);
}
// Divide prediction std/range with std of residuals to get CNR
for (unsigned int i=0; i<grpb.size(); i++) {
if (std_cnr) (*std_cnr)[i] = stdvols[i] / stdres[i];
if (range_cnr) (*range_cnr)[i] = maxvols[i] / stdres[i];
}
}
// Get the SNR of the b0s
if (b0_snr) {
if (sm.NScans(ScanType::b0) > 1) {
std::vector<unsigned int> b0_indx = sm.GetB0Indicies();
*b0_snr = b0_pmp->Predict(0) - b0_pmp->InputData(0); // N.B. Predict(i) is mean of all b0 scans
*b0_snr *= *b0_snr;
for (unsigned int i=1; i<b0_indx.size(); i++) {
NEWIMAGE::volume<float> tmp = b0_pmp->Predict(i) - b0_pmp->InputData(i);
*b0_snr += tmp*tmp;
}
*b0_snr /= static_cast<float>(b0_indx.size()-1);
*b0_snr = NEWIMAGE::sqrt(*b0_snr);
*b0_snr = b0_pmp->Predict(0) / *b0_snr;
}
else (*b0_snr) = 0.0; // Set SNR to zero if we can't estimate it
}
// Get residuals
if (residuals) {
for (unsigned int i=0; i<sm.NScans(ScanType::DWI); i++) {
(*residuals)[sm.GetDwi2GlobalIndexMapping(i)] = (dwi_pmp->InputData(i) - dwi_pmp->Predict(i)) / sm.ScaleFactor();
}
for (unsigned int i=0; i<sm.NScans(ScanType::b0); i++) {
(*residuals)[sm.Getb02GlobalIndexMapping(i)] = (b0_pmp->InputData(i) - b0_pmp->Predict(i)) / sm.ScaleFactor();
}
}
prof.EndEntry(total_key);
return;
} EddyCatch
void WriteCNRMaps(// Input
const EddyCommandLineOptions& clo,
const ECScanManager& sm,
const std::string& spatial_fname,
const std::string& range_fname,
const std::string& residual_fname) EddyTry
{
static Utilities::FSLProfiler prof("_"+string(__FILE__)+"_"+string(__func__));
double total_key = prof.StartEntry("Total");
if (spatial_fname.empty() && residual_fname.empty()) throw EddyException("EDDY::WriteCNRMaps: At least one of spatial and residual fname must be set");
// Allocate memory for the maps we are interested in
std::shared_ptr<NEWIMAGE::volume4D<float> > std_cnr;
std::shared_ptr<NEWIMAGE::volume4D<float> > range_cnr;
std::shared_ptr<NEWIMAGE::volume<float> > b0_snr;
std::shared_ptr<NEWIMAGE::volume4D<float> > residuals;
if (!spatial_fname.empty()) {
const NEWIMAGE::volume<float>& tmp = sm.Scan(0,ScanType::Any).GetIma();
std_cnr = std::make_shared<NEWIMAGE::volume4D<float> > (tmp.xsize(),tmp.ysize(),tmp.zsize(),sm.NoOfShells(ScanType::DWI));
NEWIMAGE::copybasicproperties(tmp,*std_cnr);
b0_snr = std::make_shared<NEWIMAGE::volume<float> > (tmp.xsize(),tmp.ysize(),tmp.zsize());
NEWIMAGE::copybasicproperties(tmp,*b0_snr);
}
if (!range_fname.empty()) {
const NEWIMAGE::volume<float>& tmp = sm.Scan(0,ScanType::Any).GetIma();
range_cnr = std::make_shared<NEWIMAGE::volume4D<float> > (tmp.xsize(),tmp.ysize(),tmp.zsize(),sm.NoOfShells(ScanType::DWI));
NEWIMAGE::copybasicproperties(tmp,*range_cnr);
if (b0_snr==nullptr) {
b0_snr = std::make_shared<NEWIMAGE::volume<float> > (tmp.xsize(),tmp.ysize(),tmp.zsize());
NEWIMAGE::copybasicproperties(tmp,*b0_snr);
}
}
if (!residual_fname.empty()) {
const NEWIMAGE::volume<float>& tmp = sm.Scan(0,ScanType::Any).GetIma();
residuals = std::make_shared<NEWIMAGE::volume4D<float> > (tmp.xsize(),tmp.ysize(),tmp.zsize(),sm.NScans(ScanType::Any));
NEWIMAGE::copybasicproperties(tmp,*residuals);
}
// Calculate the maps we are interested in
EDDY::CalculateCNRMaps(clo,sm,std_cnr,range_cnr,b0_snr,residuals);
// Write them out
if (!spatial_fname.empty()) {
const NEWIMAGE::volume<float>& tmp = sm.Scan(0,ScanType::Any).GetIma();
NEWIMAGE::volume4D<float> ovol(tmp.xsize(),tmp.ysize(),tmp.zsize(),sm.NoOfShells(ScanType::Any));
NEWIMAGE::copybasicproperties(tmp,ovol);
ovol[0] = *b0_snr;
for (unsigned int i=0; i<sm.NoOfShells(ScanType::DWI); i++) ovol[i+1] = (*std_cnr)[i];
NEWIMAGE::write_volume(ovol,spatial_fname);
}
if (!range_fname.empty()) {
const NEWIMAGE::volume<float>& tmp = sm.Scan(0,ScanType::Any).GetIma();
NEWIMAGE::volume4D<float> ovol(tmp.xsize(),tmp.ysize(),tmp.zsize(),sm.NoOfShells(ScanType::Any));
NEWIMAGE::copybasicproperties(tmp,ovol);
ovol[0] = *b0_snr;
for (unsigned int i=0; i<sm.NoOfShells(ScanType::DWI); i++) ovol[i+1] = (*range_cnr)[i];
NEWIMAGE::write_volume(ovol,range_fname);
}
if (!residual_fname.empty()) NEWIMAGE::write_volume(*residuals,residual_fname);
prof.EndEntry(total_key);
return;
} EddyCatch
/****************************************************************//**
*
* A global function that performs an update of the movement and EC
* parameters for a range of scans. It is a wrappper to facilitate using
* the C++11 thread library for parallelisation.
* \param[in] first_vol Index of first volume to process
* \param[in] last_vol Index one past the last volume to process
* \param[in] clo Carries information about the command line options
* that eddy was invoked with.
* \param[in] pmp A safe pointer to a DWIPredictionMaker that can be used to
* make predictions about what the scans should look like in undistorted space.
* \param[in] mask A mask specifying where predictions are valid
* \param[in] st Specifies if we should update the diffusion weighted
* images or the b=0 images.
* parameters will be checked before acceptiong them.
* \param[in] fwhm FWHM
* \param[in] iter Specifies the iteration number.
* \param[in/out] sm Collection of all scans.
* \param[out] mss The mss for the volumes after applying new parameters.
* \return None
*
********************************************************************/
void MovAndECParamUpdateWrapper(// Input
unsigned int first_vol,
unsigned int last_vol,
const EddyCommandLineOptions& clo,
std::shared_ptr<const DWIPredictionMaker> pmp,
const NEWIMAGE::volume<float>& mask,
ScanType st,
float fwhm,
unsigned int iter,
// Input/output
ECScanManager& sm,
// Output
double *mss) EddyTry
{
static std::mutex cout_mutex;
for (unsigned int s=first_vol; s<last_vol; s++) {
NEWIMAGE::volume<float> pred = pmp->Predict(s);
unsigned int global_indx = sm.GetGlobalIndex(s,st);
if (clo.DebugLevel() && clo.DebugIndicies().IsAmongIndicies(global_indx)) {
mss[s] = EddyUtils::MovAndECParamUpdate(pred,sm.GetSuscHzOffResField(s,st),sm.GetBiasField(),mask,fwhm,clo.VeryVerbose(),global_indx,iter,clo.DebugLevel(),sm.Scan(s,st));
}
else {
mss[s] = EddyUtils::MovAndECParamUpdate(pred,sm.GetSuscHzOffResField(s,st),sm.GetBiasField(),mask,fwhm,clo.VeryVerbose(),sm.Scan(s,st));
}
if (clo.VeryVerbose()) {
cout_mutex.lock();
std::cout << "Iter: " << iter << ", scan: " << global_indx << ", mss = " << mss[s] << std::endl << std::flush;
cout_mutex.unlock();
}
}
} EddyCatch
/****************************************************************//**
*
* A global function that gets an unwarped scan and loads it into a
* prediction maker. It does this for a range of scans. It is a wrappper
* to facilitate using the C++11 thread library for parallelisation.
* \param[in] first_vol Index of first volume to process
* \param[in] last_vol Index one past the last volume to process
* \param[in] clo Carries information about the command line options
* that eddy was invoked with.
* \param[in] sm Collection of all scans.
* \param[in] st Specifies if we should update the diffusion weighted
* images or the b=0 images.
* \param[in/out] pmp A safe pointer to a DWIPredictionMaker that is to
* be loaded.
* \param[in/out] mask A mask specifying where predictions are valid
* \return None
*
********************************************************************/
void SetUnwarpedScanWrapper(// Input
unsigned int first_vol,
unsigned int last_vol,
const EddyCommandLineOptions& clo,
const ECScanManager& sm,
ScanType st,
bool use_orig,
// Input/output
std::shared_ptr<DWIPredictionMaker> pmp,
NEWIMAGE::volume<float>& mask) EddyTry
{
static std::mutex cout_mutex;
static std::mutex mask_mutex;
for (unsigned int s=first_vol; s<last_vol; s++) {
if (clo.VeryVerbose()) {
cout_mutex.lock();
std::cout << " " << sm.GetGlobalIndex(s,st) << std::flush;
cout_mutex.unlock();
}
NEWIMAGE::volume<float> tmpmask = sm.Scan(s,st).GetIma();
EddyUtils::SetTrilinearInterp(tmpmask); tmpmask = 1.0;
if (use_orig) pmp->SetScan(sm.GetUnwarpedOrigScan(s,tmpmask,st),sm.Scan(s,st).GetDiffPara(clo.RotateBVecsDuringEstimation()),s);
else pmp->SetScan(sm.GetUnwarpedScan(s,tmpmask,st),sm.Scan(s,st).GetDiffPara(clo.RotateBVecsDuringEstimation()),s);
mask_mutex.lock();
mask *= tmpmask;
mask_mutex.unlock();
}
} EddyCatch
/****************************************************************//**
*
* A global function that calculates slice wise stats for a range
* of scans. It is a wrappper to facilitate using the C++11 thread
* library for parallelisation.
* \param[in] first_vol Index of first volume to process
* \param[in] last_vol Index one past the last volume to process
* \param[in] sm Collection of all scans.
* \param[in] pmp A safe pointer to a DWIPredictionMaker that is to
* be loaded.
* \param[in] mask A mask to limit the area for which to calculate stats.
* \param[in] st Specifies if we should update the diffusion weighted
* images or the b=0 images.
* \param[out] stats A vector of slice-wise stats.
* \return None
*
********************************************************************/
void GetSliceWiseStatsWrapper(// Input
unsigned int first_vol,
unsigned int last_vol,
const EddyCommandLineOptions& clo,
const ECScanManager& sm,
ScanType st,
std::shared_ptr<const DWIPredictionMaker> pmp,
const NEWIMAGE::volume<float>& mask,
unsigned int iter,
unsigned int dl,
// Output
DiffStatsVector& stats) EddyTry
{
for (unsigned int s=first_vol; s<last_vol; s++) {
NEWIMAGE::volume<float> pred = pmp->Predict(s);
stats[s] = EddyUtils::GetSliceWiseStats(clo,st,pred,sm.GetSuscHzOffResField(s,st),mask,sm,s,iter,dl);
}
} EddyCatch
/****************************************************************//**
*
* A global function that replaces outliers with predictions for a range
* of scans. It is a wrappper to facilitate using the C++11 thread
* library for parallelisation.
* \param[in] first_vol Index of first volume to process
* \param[in] last_vol Index one past the last volume to process
* \param[in] pmp A safe pointer to a DWIPredictionMaker that is to
* be loaded.
* \param[i] mask
* \param[in] st Specifies if we should update the diffusion weighted
* images or the b=0 images.
* \param[in] very_verbose Very Verbose flag
* \param[in] add_noise Decides of noise should be added to predictions.
* \param[in/out] sm Collection of all scans.
* \return None
*
********************************************************************/
void SetAsOutliersWrapper(// Input
unsigned int first_vol,
unsigned int last_vol,
std::shared_ptr<DWIPredictionMaker> pmp,
const NEWIMAGE::volume<float>& mask,
const ReplacementManager& rm,
ScanType st,
bool very_verbose,
bool add_noise,
// Input/output
ECScanManager& sm) EddyTry
{
static std::mutex cout_mutex;
for (unsigned int s=first_vol; s<last_vol; s++) {
std::vector<unsigned int> ol = rm.OutliersInScan(s);
if (ol.size()) { // If the scan has outliers
if (very_verbose) {
cout_mutex.lock();
std::cout << "Scan " << sm.GetGlobalIndex(s,st) << " has outlier slices:";
for (unsigned int i=0; i<ol.size(); i++) { std::cout << " " << ol[i]; }
std::cout << std::endl << std::flush;
cout_mutex.unlock();
}
NEWIMAGE::volume<float> pred = pmp->Predict(s,true);
if (add_noise) {
double vp = pmp->PredictionVariance(s,true);
double ve = pmp->ErrorVariance(s);
double stdev = std::sqrt(vp+ve) - std::sqrt(vp);
pred += EddyUtils::MakeNoiseIma(pred,0.0,stdev);
}
sm.Scan(s,st).SetAsOutliers(pred,sm.GetSuscHzOffResField(s,st),mask,ol);
}
}
} EddyCatch
/*
void WriteCNRMaps(// Input
const EddyCommandLineOptions& clo,
const ECScanManager& sm,
const std::string& spatial_fname,
const std::string& range_fname,
const std::string& residual_fname) EddyTry
{
if (spatial_fname == std::string("") && residual_fname == std::string("")) throw EddyException("EDDY::WriteCNRMaps: At least one of spatial and residual fname must be set");
// Load prediction maker
NEWIMAGE::volume<float> mask = sm.Scan(0,ScanType::DWI).GetIma(); EddyUtils::SetTrilinearInterp(mask); mask = 1.0; // FOV-mask in model space
#ifdef COMPILE_GPU
std::shared_ptr<DWIPredictionMaker> dwi_pmp = EddyGpuUtils::LoadPredictionMaker(clo,ScanType::DWI,sm,0,0.0,mask);
std::shared_ptr<DWIPredictionMaker> b0_pmp = EddyGpuUtils::LoadPredictionMaker(clo,ScanType::b0,sm,0,0.0,mask);
#else
std::shared_ptr<DWIPredictionMaker> dwi_pmp = EDDY::LoadPredictionMaker(clo,ScanType::DWI,sm,0,0.0,mask);
std::shared_ptr<DWIPredictionMaker> b0_pmp = EDDY::LoadPredictionMaker(clo,ScanType::b0,sm,0,0.0,mask);
#endif
if (sm.IsShelled()) {
if (spatial_fname != std::string("") || range_fname != std::string("")) { // Calculate and write shell-wise CNR maps
std::vector<double> grpb;
std::vector<std::vector<unsigned int> > dwi_indx = sm.GetShellIndicies(grpb); // Global indicies of dwi scans
std::vector<NEWIMAGE::volume<float> > mvols(grpb.size());
// Calculate mean volumes of predictions
for (unsigned int i=0; i<grpb.size(); i++) {
mvols[i] = dwi_pmp->Predict(sm.GetGlobal2DWIIndexMapping(dwi_indx[i][0]));
for (unsigned int j=1; j<dwi_indx[i].size(); j++) {
mvols[i] += dwi_pmp->Predict(sm.GetGlobal2DWIIndexMapping(dwi_indx[i][j]));
}
mvols[i] /= float(dwi_indx[i].size());
}
// Calculate standard deviation volumes of predictions
std::vector<NEWIMAGE::volume<float> > stdvols(grpb.size());
for (unsigned int i=0; i<grpb.size(); i++) {
NEWIMAGE::volume<float> tmp = dwi_pmp->Predict(sm.GetGlobal2DWIIndexMapping(dwi_indx[i][0])) - mvols[i];
stdvols[i] = tmp*tmp;
for (unsigned int j=1; j<dwi_indx[i].size(); j++) {
tmp = dwi_pmp->Predict(sm.GetGlobal2DWIIndexMapping(dwi_indx[i][j])) - mvols[i];
stdvols[i] += tmp*tmp;
}
stdvols[i] /= float(dwi_indx[i].size()-1);
stdvols[i] = NEWIMAGE::sqrt(stdvols[i]);
}
// Calculate range (min--max) volumes of predictions to make dHCP data seem like it has decent CNR
std::vector<NEWIMAGE::volume<float> > minvols(grpb.size());
std::vector<NEWIMAGE::volume<float> > maxvols(grpb.size());
for (unsigned int i=0; i<grpb.size(); i++) {
minvols[i] = dwi_pmp->Predict(sm.GetGlobal2DWIIndexMapping(dwi_indx[i][0]));
maxvols[i] = minvols[i];
for (unsigned int j=1; j<dwi_indx[i].size(); j++) {
minvols[i] = NEWIMAGE::min(minvols[i],dwi_pmp->Predict(sm.GetGlobal2DWIIndexMapping(dwi_indx[i][j])));
maxvols[i] = NEWIMAGE::max(maxvols[i],dwi_pmp->Predict(sm.GetGlobal2DWIIndexMapping(dwi_indx[i][j])));
}
maxvols[i] -= minvols[i]; // Maxvols now contain the range rather than the max
}
// Calculate standard deviation of residuals
std::vector<NEWIMAGE::volume<float> > stdres(grpb.size());
for (unsigned int i=0; i<grpb.size(); i++) {
NEWIMAGE::volume<float> tmp = dwi_pmp->Predict(sm.GetGlobal2DWIIndexMapping(dwi_indx[i][0])) - dwi_pmp->InputData(sm.GetGlobal2DWIIndexMapping(dwi_indx[i][0]));
stdres[i] = tmp*tmp;
for (unsigned int j=1; j<dwi_indx[i].size(); j++) {
tmp = dwi_pmp->Predict(sm.GetGlobal2DWIIndexMapping(dwi_indx[i][j])) - dwi_pmp->InputData(sm.GetGlobal2DWIIndexMapping(dwi_indx[i][j]));
stdres[i] += tmp*tmp;
}
stdres[i] /= float(dwi_indx[i].size()-1);
stdres[i] = NEWIMAGE::sqrt(stdres[i]);
}
// Divide std and range of predictions with std of residuals to get CNR
for (unsigned int i=0; i<grpb.size(); i++) {
maxvols[i] /= stdres[i];
stdvols[i] /= stdres[i];
}
// Get the SNR (since CNR is zero) of the b0s
std::vector<unsigned int> b0_indx = sm.GetB0Indicies();
NEWIMAGE::volume<float> b0_SNR = b0_pmp->Predict(0) - b0_pmp->InputData(0);
if (b0_indx.size() > 1) {
b0_SNR *= b0_SNR;
for (unsigned int i=1; i<b0_indx.size(); i++) {
NEWIMAGE::volume<float> tmp = b0_pmp->Predict(i) - b0_pmp->InputData(i);
b0_SNR += tmp*tmp;
}
b0_SNR /= float(b0_indx.size()-1);
b0_SNR = NEWIMAGE::sqrt(b0_SNR);
b0_SNR = b0_pmp->Predict(0) /= b0_SNR;
}
else b0_SNR = 0.0; // Set it to zero if we can't assess it.
// Put SNR and CNR maps together into 4D file with spatial CNR
NEWIMAGE::volume4D<float> spat_cnr(stdvols[0].xsize(),stdvols[0].ysize(),stdvols[0].zsize(),stdvols.size()+1);
NEWIMAGE::copybasicproperties(stdvols[0],spat_cnr);
spat_cnr[0] = b0_SNR;
if (spatial_fname != std::string("")) {
for (unsigned int i=0; i<stdvols.size(); i++) spat_cnr[i+1] = stdvols[i];
NEWIMAGE::write_volume(spat_cnr,spatial_fname);
}
// Put SNR and range maps together into 4D file with spatial "range-CNR"
if (range_fname != std::string("")) {
for (unsigned int i=0; i<maxvols.size(); i++) spat_cnr[i+1] = maxvols[i];
NEWIMAGE::write_volume(spat_cnr,range_fname);
}
}
if (residual_fname != std::string("")) { // Calculate and write maps of residuals for all (b0 and DWI) scans
NEWIMAGE::volume4D<float> residuals(dwi_pmp->InputData(0).xsize(),dwi_pmp->InputData(0).ysize(),dwi_pmp->InputData(0).zsize(),sm.NScans(ScanType::Any));
NEWIMAGE::copybasicproperties(dwi_pmp->InputData(0),residuals);
for (unsigned int i=0; i<sm.NScans(ScanType::DWI); i++) {
residuals[sm.GetDwi2GlobalIndexMapping(i)] = dwi_pmp->InputData(i) - dwi_pmp->Predict(i);
}
for (unsigned int i=0; i<sm.NScans(ScanType::b0); i++) {
residuals[sm.Getb02GlobalIndexMapping(i)] = b0_pmp->InputData(i) - b0_pmp->Predict(i);
}
NEWIMAGE::write_volume(residuals,residual_fname);
}
}
else {
throw EddyException("WriteCNRMaps: Cannot calculate CNR for non-shelled data.");
}
} EddyCatch
*/
void Diagnostics(// Input
const EddyCommandLineOptions& clo,
unsigned int iter,
ScanType st,
const ECScanManager& sm,
const double *mss_tmp,
const DiffStatsVector& stats,
const ReplacementManager& rm,
// Output
NEWMAT::Matrix& mss,
NEWMAT::Matrix& phist) EddyTry
{
if (clo.Verbose()) {
double tss=0.0;
for (unsigned int s=0; s<sm.NScans(st); s++) tss+=mss_tmp[s];
cout << "Iter: " << iter << ", Total mss = " << tss/sm.NScans(st) << endl;
}
for (unsigned int s=0; s<sm.NScans(st); s++) {
mss(iter+1,s+1) = mss_tmp[s];
phist.SubMatrix(iter+1,iter+1,s*sm.Scan(0,st).NParam()+1,(s+1)*sm.Scan(0,st).NParam()) = sm.Scan(s,st).GetParams().t();
}
if (clo.WriteSliceStats()) {
char istring[256];
if (st==EDDY::ScanType::DWI) sprintf(istring,"%s.EddyDwiSliceStatsIteration%02d",clo.IOutFname().c_str(),iter);
else sprintf(istring,"%s.Eddyb0SliceStatsIteration%02d",clo.IOutFname().c_str(),iter);
stats.Write(string(istring));
rm.DumpOutlierMaps(string(istring));
}
} EddyCatch
void AddRotation(ECScanManager& sm,
const NEWMAT::ColumnVector& rp) EddyTry
{
for (unsigned int i=0; i<sm.NScans(); i++) {
NEWMAT::ColumnVector mp = sm.Scan(i).GetParams(ParametersType::Movement);
mp(4) += rp(1); mp(5) += rp(2); mp(6) += rp(3);
sm.Scan(i).SetParams(mp,ParametersType::Movement);
}
} EddyCatch
void PrintMIValues(const EddyCommandLineOptions& clo,
const ECScanManager& sm,
const std::string& fname,
bool write_planes) EddyTry
{
std::vector<std::string> dir(6);
dir[0]="xt"; dir[1]="yt"; dir[2]="zt";
dir[3]="xr"; dir[4]="yr"; dir[5]="zr";
// First write 1D profiles along main directions
for (unsigned int i=0; i<6; i++) {
std::vector<unsigned int> n(6,1); n[i] = 100;
std::vector<double> first(6,0.0); first[i] = -2.5;
std::vector<double> last(6,0.0); last[i] = 2.5;
if (clo.VeryVerbose()) cout << "Writing MI values for direction " << i << endl;
PEASUtils::WritePostEddyBetweenShellMIValues(clo,sm,n,first,last,fname+"_"+dir[i]);
}
// Write 2D planes if requested
for (unsigned int i=0; i<6; i++) {
for (unsigned int j=i+1; j<6; j++) {
std::vector<unsigned int> n(6,1); n[i] = 20; n[j] = 20;
std::vector<double> first(6,0.0); first[i] = -1.0; first[j] = -1.0;
std::vector<double> last(6,0.0); last[i] = 1.0; last[j] = 1.0;
if (clo.VeryVerbose()) cout << "Writing MI values for plane " << i << "-" << j << endl;
PEASUtils::WritePostEddyBetweenShellMIValues(clo,sm,n,first,last,fname+"_"+dir[i]+"_"+dir[j]);
}
}
} EddyCatch
void json_out_struct::Write(const std::string& fname) const EddyTry
{
nlohmann::json json_file;
json_file["Shell indicies"] = this->shell_indicies;
json_file["Outlier report"] = this->outlier_report;
json_file["Outlier maps"] = this->outlier_maps;
json_file["Parameters"] = this->parameters;
json_file["Movement over time"] = this->movement_over_time;
json_file["Movement RMS"] = this->movement_rms;
json_file["Restricted movement RMS"] = this->restricted_movement_rms;
json_file["Long time-constant EC parameters"] = this->long_ec_parameters;
try {
std::ofstream ofile(fname);
ofile << std::setw(4) << json_file << std::endl;
ofile.close();
}
catch(const std::exception& e) {
throw EddyException("json_out_struct::Write: Exception thrown with message " + std::string(e.what()));
}
} EddyCatch
void WritePMDebugInfo(std::shared_ptr<DWIPredictionMaker> pmp,
const EddyCommandLineOptions& clo,
EDDY::ScanType st,
unsigned int iter,
const std::string& arch,
const std::string& when) EddyTry
{
std::string prefix("EDDY_DEBUG_");
prefix += arch;
if (st==ScanType::b0) prefix += "b0";
else if (st==ScanType::DWI) prefix += "DWI";
else if (st==ScanType::fMRI) prefix += "fMRI";
if (clo.DebugLevel() > 2) {
char fname[256];
sprintf(fname,"%s_K_Mat_Data_%s_%02d",prefix.c_str(),when.c_str(),iter);
pmp->WriteMetaData(fname);
}
if (clo.DebugLevel() > 3) {
char fname[256];
sprintf(fname,"%s_K_Mat_Ima_%s_%02d",prefix.c_str(),when.c_str(),iter);
pmp->WriteImageData(fname);
}
} EddyCatch
} // End namespace EDDY
/*! \mainpage
* Here goes a description of the eddy project
*/
eddy.h 0 → 100644
View file @ b9e18ed6
/////////////////////////////////////////////////////////////////////
///
/// \file eddy.h
/// \brief Contains declarations of some very high level functions for eddy
///
/// This file contains declarations for some very high level functions
/// that are called in eddy.cpp.
///
/// \author Jesper Andersson
/// \version 1.0b, Nov., 2012.
/// \Copyright (C) 2012 University of Oxford
///
/////////////////////////////////////////////////////////////////////
#ifndef eddy_h
#define eddy_h
#include <cstdlib>
#include <string>
#include <vector>
#include <cmath>
#include <memory>
#include "armawrap/newmat.h"
#include "EddyHelperClasses.h"
#include "ECScanClasses.h"
#include "DiffusionGP.h"
#include "b0Predictor.h"
#include "EddyUtils.h"
#include "EddyCommandLineOptions.h"
namespace EDDY {
/// A very high-level global function that registers all the scans (b0 and dwis) in sm using a volume-to-volume model.
ReplacementManager *DoVolumeToVolumeRegistration(// Input
const EddyCommandLineOptions& clo,
// Input/Output
ECScanManager& sm);
/// A very high-level global function that estimates long time-constant EC and (optionally) and re-estimates movement and EC.
ReplacementManager *EstimateLongEC(// Input
const EddyCommandLineOptions& clo,
// Input/Output
ECScanManager& sm,
ReplacementManager *rm);
/// A very high-level global function that registers all the scans (b0 and dwis) in sm using a slice-to-volume model.
ReplacementManager *DoSliceToVolumeRegistration(// Input
const EddyCommandLineOptions& clo,
unsigned int oi, // Order index
bool dol, // Detect outliers?
// Input/Output
ECScanManager& sm,
ReplacementManager *dwi_rm);
/// A very high-level global function that estimates a bias field
void EstimateBiasField(// Input
const EddyCommandLineOptions& clo,
double ksp,
double lambda,
// Input/output
ECScanManager& sm);
/// Global function that registers a set of scans together
ReplacementManager *Register(// Input
const EddyCommandLineOptions& clo,
ScanType st,
unsigned int niter,
const std::vector<float>& fwhm,
SecondLevelECModelType slm,
bool dol,
// Input/Output
ECScanManager& sm,
ReplacementManager *rm,
// Output
NEWMAT::Matrix& msshist,
NEWMAT::Matrix& phist);
/// Global function that performs final check for outliers without error-variance fudging.
ReplacementManager *FinalOLCheck(// Input
const EddyCommandLineOptions& clo,
// Input/output
ReplacementManager *rm,
ECScanManager& sm);
/// Global function that detect outlier slices and replaces them by their expectation
DiffStatsVector DetectAndReplaceOutliers(// Input
const EddyCommandLineOptions& clo,
ScanType st,
// Input/Output
ECScanManager& sm,
ReplacementManager& rm);
/// Global function that Loads up the prediction maker with unwarped scans
std::shared_ptr<DWIPredictionMaker> LoadPredictionMaker(// Input
const EddyCommandLineOptions& clo,
ScanType st,
const ECScanManager& sm,
unsigned int iter,
float fwhm,
// Output
NEWIMAGE::volume<float>& mask,
// Optional input
bool use_orig=false);
/*
/// Global function that replaces selected (by rm) volumes in prediction maker
void UpdatePredictionMaker(// Input
const EddyCommandLineOptions& clo,
ScanType st,
const ECScanManager& sm,
const ReplacementManager& rm,
const NEWIMAGE::volume<float>& mask,
// Input/Output
std::shared_ptr<DWIPredictionMaker> pmp);
*/
/// Looks for outlier slices
DiffStatsVector DetectOutliers(// Input
const EddyCommandLineOptions& clo,
ScanType st,
std::shared_ptr<const DWIPredictionMaker> pmp,
const NEWIMAGE::volume<float>& mask,
const ECScanManager& sm,
// Input/Output
ReplacementManager& rm);
DiffStatsVector DetectOutliers(// Input
const EddyCommandLineOptions& clo,
ScanType st,
std::shared_ptr<const DWIPredictionMaker> pmp,
const NEWIMAGE::volume<float>& mask,
const ECScanManager& sm,
unsigned int iter,
unsigned int dl, // Debug Level
// Input/Output
ReplacementManager& rm);
DiffStatsVector detect_outliers(// Input
const EddyCommandLineOptions& clo,
ScanType st,
std::shared_ptr<const DWIPredictionMaker> pmp,
const NEWIMAGE::volume<float>& mask,
const ECScanManager& sm,
unsigned int iter,
unsigned int dl, // Debug Level
// Input/Output
ReplacementManager& rm);
/// Replaces outlier slices with their predictions
void ReplaceOutliers(// Input
const EddyCommandLineOptions& clo,
ScanType st,
std::shared_ptr<DWIPredictionMaker> pmp,
const NEWIMAGE::volume<float>& mask,
const ReplacementManager& rm,
bool add_noise,
// Input/Output
ECScanManager& sm);
/// Get predictions to help with slice-to-vol resampling
std::vector<double> GetPredictionsForResampling(// Input
const EddyCommandLineOptions& clo,
ScanType st,
const ECScanManager& sm,
// Output
NEWIMAGE::volume4D<float>& pred);
void GetScatterBrainPredictions(// Input
const EddyCommandLineOptions& clo,
ScanType st,
ECScanManager& sm,
const std::vector<double>& hypar,
// Output
NEWIMAGE::volume4D<float>& pred,
// Optional input
bool vwbvrot=false);
/// Calculate maps of CNR and SNR
void CalculateCNRMaps(// Input
const EddyCommandLineOptions& clo,
const ECScanManager& sm,
// Output
std::shared_ptr<NEWIMAGE::volume4D<float> > std_cnr,
std::shared_ptr<NEWIMAGE::volume4D<float> > range_cnr,
std::shared_ptr<NEWIMAGE::volume<float> > b0_snr,
std::shared_ptr<NEWIMAGE::volume4D<float> > residuals);
/// Write maps of CNR and SNR
void WriteCNRMaps(// Input
const EddyCommandLineOptions& clo,
const ECScanManager& sm,
const std::string& spatial_fname,
const std::string& range_fname,
const std::string& temporal_fname);
/// A global function that performs an update of the movement and EC parameters for a range of
/// scans. It is a wrappper to facilitate using the C++11 thread library for parallelisation.
void MovAndECParamUpdateWrapper(// Input
unsigned int first_vol,
unsigned int last_vol,
const EddyCommandLineOptions& clo,
std::shared_ptr<const DWIPredictionMaker> pmp,
const NEWIMAGE::volume<float>& mask,
ScanType st,
float fwhm,
unsigned int iter,
// Input/output
ECScanManager& sm,
// Output
double *mss);
/// A global function that unwarps a range of scans and inserts them into a prediction
/// maker. It is a wrappper to facilitate using the C++11 thread library for parallelisation.
void SetUnwarpedScanWrapper(// Input
unsigned int first_vol,
unsigned int last_vol,
const EddyCommandLineOptions& clo,
const ECScanManager& sm,
ScanType st,
bool use_orig,
// Input/output
std::shared_ptr<DWIPredictionMaker> pmp,
NEWIMAGE::volume<float>& mask);
/// A global function that calculates slicestats and replaces outliers for a range of scans.
/// It is a wrappper to facilitate using the C++11 thread library for parallelisation.
void GetSliceWiseStatsWrapper(// Input
unsigned int first_vol,
unsigned int last_vol,
const EddyCommandLineOptions& clo,
const ECScanManager& sm,
ScanType st,
std::shared_ptr<const DWIPredictionMaker> pmp,
const NEWIMAGE::volume<float>& mask,
unsigned int iter,
unsigned int dl,
// Output
DiffStatsVector& stats);
/// A global function that replaces outliers with predictions for a range of scans.
/// It is a wrappper to facilitate using the C++11 thread library for parallelisation.
void SetAsOutliersWrapper(// Input
unsigned int first_vol,
unsigned int last_vol,
std::shared_ptr<DWIPredictionMaker> pmp,
const NEWIMAGE::volume<float>& mask,
const ReplacementManager& rm,
ScanType st,
bool very_verbose,
bool add_noise,
// Input/output
ECScanManager& sm);
/// Global function that Generates diagnostic information for subsequent analysis
void Diagnostics(// Input
const EddyCommandLineOptions& clo,
unsigned int iter,
ScanType st,
const ECScanManager& sm,
const double *mss_tmp,
const DiffStatsVector& stats,
const ReplacementManager& rm,
// Output
NEWMAT::Matrix& mss,
NEWMAT::Matrix& phist);
/// Global function that adds rotation to all volumes.
void AddRotation(ECScanManager& sm,
const NEWMAT::ColumnVector& rp);
/// Global function that prints out (for debugging) MI values between shells
void PrintMIValues(const EddyCommandLineOptions& clo,
const ECScanManager& sm,
const std::string& fname,
bool write_planes);
/// Global function that writes prediction-maker debug information
void WritePMDebugInfo(std::shared_ptr<DWIPredictionMaker> pmp,
const EddyCommandLineOptions& clo,
EDDY::ScanType st,
unsigned int iter,
const std::string& arch,
const std::string& when);
} // End namespace EDDY
#endif // End #ifndef eddy_h
eddy_cpu 0 → 100755
View file @ b9e18ed6
File added
eddy_cuda 0 → 100755
View file @ b9e18ed6
File added
fmriPredictor.cpp 0 → 100644
View file @ b9e18ed6
/*! \file fmriPredictor.cpp
\brief Contains definitions for class for making Gaussian process based predictions about fMRI data.
\author Jesper Andersson
\version 1.0b, Feb., 2022.
*/
// Definitions of class to make Gaussian-Process
// based predictions about diffusion data.
//
// fmriPredictor.cpp
//
// Jesper Andersson, FMRIB Image Analysis Group
//
// Copyright (C) 2022 University of Oxford
//
#include <cstdlib>
#include <string>
#include <vector>
#include <cmath>
#include "armawrap/newmat.h"
#include "newimage/newimageall.h"
#include "miscmaths/miscmaths.h"
#include "EddyHelperClasses.h"
#include "EddyUtils.h"
#include "KMatrix.h"
#include "HyParEstimator.h"
#include "fmriPredictor.h"
using namespace EDDY;
NEWIMAGE::volume<float> fmriPredictor::Predict(unsigned int indx,
bool exclude) const EddyTry
{
if (!IsPopulated()) throw EddyException("fmriPredictor::Predict:const: Not yet fully populated");
if (!IsValid()) throw EddyException("fmriPredictor::Predict:const: Not yet ready for predictions");
arma::rowvec pv = _Kmats[_glist[indx]._sess]->PredVec(_glist[indx]._sindx,exclude); // Calls const version
NEWIMAGE::volume<float> pi = *(_slist[0].SPtr(0)); pi = 0.0;
#ifdef COMPILE_GPU
predict_image_gpu(indx,exclude,pv,pi);
#else
predict_image_cpu(indx,exclude,pv,pi);
#endif
return(pi);
} EddyCatch
NEWIMAGE::volume<float> fmriPredictor::Predict(unsigned int indx,
bool exclude) EddyTry
{
if (!IsPopulated()) throw EddyException("fmriPredictor::Predict:non-const: Not yet fully populated");
if (!IsValid()) throw EddyException("fmriPredictor::Predict:non-const: Not yet ready for predictions");
arma::rowvec pv = _Kmats[_glist[indx]._sess]->PredVec(_glist[indx]._sindx,exclude); // Calls non-const version
NEWIMAGE::volume<float> pi = *(_slist[0].SPtr(0)); pi = 0.0;
#ifdef COMPILE_GPU
predict_image_gpu(indx,exclude,pv,pi);
#else
predict_image_cpu(indx,exclude,pv,pi);
#endif
return(pi);
} EddyCatch
NEWIMAGE::volume<float> fmriPredictor::PredictCPU(unsigned int indx,
bool exclude) EddyTry
{
if (!IsPopulated()) throw EddyException("fmriPredictor::Predict:non-const: Not yet fully populated");
if (!IsValid()) throw EddyException("fmriPredictor::Predict:non-const: Not yet ready for predictions");
arma::rowvec pv = _Kmats[_glist[indx]._sess]->PredVec(_glist[indx]._sindx,exclude); // Calls non-const version
NEWIMAGE::volume<float> pi = *(_slist[0].SPtr(0)); pi = 0.0;
predict_image_cpu(indx,exclude,pv,pi);
return(pi);
} EddyCatch
std::vector<NEWIMAGE::volume<float> > fmriPredictor::Predict(const std::vector<unsigned int>& indicies,
bool exclude) EddyTry
{
if (!IsPopulated()) throw EddyException("fmriPredictor::Predict: Not yet fully populated");
if (!IsValid()) throw EddyException("fmriPredictor::Predict: Not yet ready for predictions");
std::vector<NEWIMAGE::volume<float> > pi(indicies.size());
std::vector<arma::rowvec> pvecs(indicies.size());
for (unsigned int i=0; i<indicies.size(); i++) pvecs[i] = _Kmats[_glist[indicies[i]]._sess]->PredVec(_glist[indicies[i]]._sindx,exclude);
#ifdef COMPILE_GPU
predict_images_gpu(indicies,exclude,pvecs,pi);
#else
predict_images_cpu(indicies,exclude,pvecs,pi);
#endif
return(pi);
} EddyCatch
NEWIMAGE::volume<float> fmriPredictor::InputData(unsigned int indx) const EddyTry
{
if (!IsPopulated()) throw EddyException("fmriPredictor::InputData: Not yet fully populated");
if (indx >= _glist.size()) throw EddyException("DiffusionGP::InputData: indx out of range");
return(*(_slist[_glist[indx]._sess].SPtr(_glist[indx]._sindx)) + *(_mptrs[_glist[indx]._sess]));
} EddyCatch
std::vector<NEWIMAGE::volume<float> > fmriPredictor::InputData(const std::vector<unsigned int>& indx) const EddyTry
{
if (!IsPopulated()) throw EddyException("fmriPredictor::InputData: Not yet fully populated");
for (unsigned int i=0; i<indx.size(); i++) if (indx[i] >= _glist.size()) throw EddyException("fmriPredictor::InputData: indx out of range");
std::vector<NEWIMAGE::volume<float> > rval(indx.size());
for (unsigned int i=0; i<indx.size(); i++) rval[i] = *(_slist[_glist[indx[i]]._sess].SPtr(_glist[indx[i]]._sindx)) + *(_mptrs[_glist[indx[i]]._sess]);
return(rval);
} EddyCatch
double fmriPredictor::PredictionVariance(unsigned int indx,
bool exclude) EddyTry
{
if (!IsPopulated()) throw EddyException("fmriPredictor::PredictionVariance:const: Not yet fully populated");
if (!IsValid()) throw EddyException("fmriPredictor::PredictionVariance:const: Not yet ready for predictions");
double pv = _Kmats[_glist[indx]._sess]->PredVar(_glist[indx]._sindx,exclude);
return(pv);
} EddyCatch
double fmriPredictor::ErrorVariance(unsigned int indx) const EddyTry
{
if (!IsPopulated()) throw EddyException("fmriPredictor::ErrorVariance:const: Not yet fully populated");
if (!IsValid()) throw EddyException("fmriPredictor::ErrorVariance:const: Not yet ready for predictions");
double ev = _Kmats[_glist[indx]._sess]->ErrVar(_glist[indx]._sindx);
return(ev);
} EddyCatch
void fmriPredictor::SetNoOfScans(unsigned int n) EddyTry
{
if (n == _glist.size()) return; // No change
else if (n > _glist.size()) { // If increasing size
std::lock_guard<std::mutex> lg(_set_mut);
_glist.resize(n); // New elements populated according to default constructor
_lak = false;
for (unsigned int i; i<_slist.size(); i++) _Kmats[i]->Reset();
}
else { // Decreasing size not allowed
throw EddyException("fmriPredictor::SetNoOfScans: Decreasing size not allowed");
}
return;
} EddyCatch
void fmriPredictor::SetScan(const NEWIMAGE::volume<float>& scan,
const DiffPara& dp,
unsigned int indx,
unsigned int sess) EddyTry
{
std::lock_guard<std::mutex> lg(_set_mut);
// First just sanity checks
if (indx >= _glist.size()) throw EddyException("fmriPredictor::SetScan: Invalid image index");
if (_tr < 0.0) _tr = dp.TR();
else if (std::abs(_tr - dp.TR()) > 1e-6) throw EddyException("fmriPredictor::SetScan: You cannot mix scans with different repetition times");
if (_glist.size() && !NEWIMAGE::samesize(*_slist[_glist[0]._sess].SPtr(0),scan)) throw EddyException("fmriPredictor::SetScan: Wrong image dimension");
if (_glist[indx]._sess >= 0 && _glist[indx]._sess != sess) throw EddyException("fmriPredictor::SetScan: You cannot change session of a scan");
if (_glist[indx]._sess >= 0 && _slist[_glist[indx]._sess].Gndx(_glist[indx]._sindx) != static_cast<int>(indx)) throw EddyException("fmriPredictor::SetScan: Lists are inconsistent");
// Next we do the job
if (_glist[indx]._sess >= 0) { // If the same index has already been loaded before
_slist[_glist[indx]._sess].SPtr(_glist[indx]._sindx) = std::make_shared<NEWIMAGE::volume<float> >(scan);
}
else { // If this is the first time this index is loaded
_glist[indx]._sess = sess;
_glist[indx]._sindx = _slist[sess].Size();
_slist[sess].PushBack(scan,indx);
}
_lak = false;
} EddyCatch
void fmriPredictor::EvaluateModel(const NEWIMAGE::volume<float>& mask, float fwhm, bool verbose) EddyTry
{
// The first thing we want to do is to "clean" up the lists.
if (!lists_are_kosher()) throw EddyException("fmriPredictor::lists_are_kosher: Lists are inconsistent");
// Next make one K-matrix per session
_Kmats.resize(_slist.size());
for (unsigned int i=1; i<_slist.size(); i++) {
_Kmats[i] = _Kmats[0]->Clone();
}
// Populate K-matrices with distances. Note that we have shoe-horned it into a "diffusion syntax"
for (unsigned int i=0; i<_slist.size(); i++) {
std::vector<DiffPara> dpars(_slist[i].Size(),DiffPara(-100,_tr));
_Kmats[i]->SetDiffusionPar(dpars);
}
mean_correct(); // Mean correct (on a per session basis)
// Next we need to select data for the estimation of hyper-parameters.
// If all sessions have the same number of scans we can mix-and-match data
// from different sessions. If not we pick a session at random, but with
// probability weighted by the number of scans in the session.
if (same_no_of_scans_in_all_sessions()) { // Mix-and-match
DataSelector ds = DataSelector(_slist[0].SPtr_list(),mask,(_hpe->GetNVox()/_slist.size())+1,FourthDimension(0),fwhm,_hpe->RndInit());
for (unsigned int i=1; i<_slist.size(); i++) {
ds.ConcatToMe(DataSelector(_slist[i].SPtr_list(),mask,(_hpe->GetNVox()/_slist.size())+1,FourthDimension(i),fwhm,_hpe->RndInit()));
}
_hpe->SetData(ds.GetData());
_hpe->Estimate(_Kmats[0],verbose);
for (unsigned int i=0; i<_slist.size(); i++) {
_Kmats[i]->SetHyperPar(_hpe->GetHyperParameters());
_Kmats[i]->CalculateInvK();
}
}
else { // Pick random session
// Draw random number 0--no_of_scans-1
int seed = (_hpe->RndInit()) ? _hpe->RndInit() : static_cast<int>(time(NULL));
std::mt19937 gen(seed);
std::uniform_int_distribution<> distr(0,static_cast<int>(no_of_scans())-1);
int rndnr = distr(gen);
// Determine which session this corresponds to
int cumscans = 0;
unsigned int si;
for (si=0; si<_slist.size(); si++) {
cumscans += _slist[si].Size();
if (cumscans >= rndnr) break;
}
DataSelector ds = DataSelector(_slist[si].SPtr_list(),mask,_hpe->GetNVox(),FourthDimension(si),fwhm,_hpe->RndInit());
_hpe->SetData(ds.GetData());
_hpe->Estimate(_Kmats[si],verbose);
for (unsigned int i=0; i<_slist.size(); i++) {
_Kmats[i]->SetHyperPar(_hpe->GetHyperParameters());
_Kmats[i]->CalculateInvK();
}
}
return;
} EddyCatch
void fmriPredictor::WriteImageData(const std::string& fname) const EddyTry
{
char ofname[256];
if (!IsPopulated() || !_lak) throw EddyException("fmriPredictor::WriteImageData: Not yet fully populated");
// For practical reasons the volumes are written individually
for (unsigned int i=0; i<_glist.size(); i++) {
sprintf(ofname,"%s_%03d_%02d_%03d",fname.c_str(),i,_glist[i]._sess,_glist[i]._sindx);
NEWIMAGE::write_volume(*(_slist[_glist[i]._sess].SPtr(_glist[i]._sindx)),ofname);
}
for (unsigned int i=0; i<_mptrs.size(); i++) {
sprintf(ofname,"%s_mean_%02d",fname.c_str(),i);
NEWIMAGE::write_volume(*(_mptrs[i]),ofname);
}
} EddyCatch
void fmriPredictor::WriteMetaData(const std::string& fname) const EddyTry
{
if (!IsPopulated() || !_lak) throw EddyException("fmriPredictor::WriteMetaData: Not yet fully populated");
char ofname[256];
for (unsigned int i=0; i<_slist.size(); i++) {
sprintf(ofname,"%s_%02d",fname.c_str(),i);
_Kmats[i]->Write(ofname);
}
} EddyCatch
void fmriPredictor::mean_correct() EddyTry
{
NEWIMAGE::volume<float> mean = *_slist[_glist[0]._sess].SPtr(_glist[0]._sindx);
for (unsigned int li=0; li<_slist.size(); li++) {
mean = *_slist[li].SPtr(0);
for (unsigned int i=1; i<_slist[li].Size(); i++) mean += *_slist[li].SPtr(i);
mean /= static_cast<float>(_slist[li].Size());
for (unsigned int i=0; i<_slist[li].Size(); i++) *_slist[li].SPtr(i) -= mean;
}
} EddyCatch
bool fmriPredictor::is_populated() const EddyTry
{
for (unsigned int i=0; i<_glist.size(); i++) if (_glist[i]._sess < 0) return(false);
return(true);
} EddyCatch
bool fmriPredictor::same_no_of_scans_in_all_sessions() const EddyTry
{
for (unsigned int i=1; i<_slist.size(); i++) if (_slist[i].Size() != _slist[0].Size()) return(false);
return(true);
} EddyCatch
/****************************************************************//**
*
* This ensures that the entries in the _slist are sorted in ascending
* order of "global index". But when reorganising the _slist we also
* need to make sure that the entries in the _glist still points to
* the right _slist entries.
*
********************************************************************/
bool fmriPredictor::lists_are_kosher() EddyTry
{
if (!_lak) {
// First sort the entries in the _slists with the _gindx as key
for (unsigned int li=0; li<_slist.size(); li++) _slist[li].SortByGndx();
// Then make sure that _glist and _slist are consistent
for (unsigned int li=0; li<_slist.size(); li++) {
for (unsigned int i=0; i<_slist[li].Size(); i++) {
if (_glist[_slist[li].Gndx(i)]._sess != li) return(false);
_glist[_slist[li].Gndx(i)]._sindx = i;
}
}
// Then do a little sanity check of _glist
for (unsigned int i=0; i<_glist.size(); i++) {
if (_glist[i]._sess >= _slist.size() || _glist[i]._sindx >= _slist[_glist[i]._sess].Size()) return(false);
}
// And finally check that _glist don't have duplicates
if (glist_has_duplicates()) return(false);
}
return(true);
} EddyCatch
bool fmriPredictor::glist_has_duplicates() const EddyTry
{
for (unsigned int i=0; i<_glist.size()-1; i++) {
for (unsigned int j=i+1; j<_glist.size(); j++) {
if (_glist[i]._sess == _glist[j]._sess && _glist[i]._sindx == _glist[j]._sindx) return(true);
}
}
return(false);
} EddyCatch
void fmriPredictor::predict_image_cpu(// Input
unsigned int indx,
bool exclude,
const arma::rowvec& pv,
// Output
NEWIMAGE::volume<float>& pi) const EddyTry
{
unsigned int ys = (exclude) ? _slist[_glist[indx]._sess].Size()-1 : _slist[_glist[indx]._sess].Size();
arma::colvec y(ys);
pi = *_mptrs[_glist[indx]._sess];
for (int k=0; k<pi.zsize(); k++) {
for (int j=0; j<pi.ysize(); j++) {
for (int i=0; i<pi.xsize(); i++) {
if (get_y(i,j,k,indx,exclude,y)) pi(i,j,k) += static_cast<float>(arma::as_scalar(pv*y));
else pi(i,j,k) = 0.0;
}
}
}
} EddyCatch
void fmriPredictor::predict_images_cpu(// Input
const std::vector<unsigned int>& indicies,
bool exclude,
const std::vector<arma::rowvec>& pvecs,
// Output
std::vector<NEWIMAGE::volume<float> >& pi) const EddyTry
{
if (indicies.size() != pvecs.size() || indicies.size() != pi.size()) {
throw EDDY::EddyException("fmriPredictor::predict_images_cpu: mismatch among indicies, pvecs and pi");
}
for (unsigned int i=0; i<indicies.size(); i++) predict_image_cpu(indicies[i],exclude,pvecs[i],pi[i]);
return;
} EddyCatch
bool fmriPredictor::get_y(// Input
unsigned int i, unsigned int j, unsigned int k, unsigned int indx, bool exclude,
// Output
arma::colvec& y) const EddyTry
{
const std::vector<std::shared_ptr<NEWIMAGE::volume<float> > >& sptrl = _slist[_glist[indx]._sess].SPtr_list();
unsigned int sindx = _glist[indx]._sindx;
for (unsigned int t=0, tt=0; t<sptrl.size(); t++) {
if (!exclude || t!=sindx) {
if (!(*sptrl[t])(i,j,k)) return(false);
else y(tt++) = (*sptrl[t])(i,j,k);
}
}
return(true);
} EddyCatch
void fmriPredictor::ptr_index_list::SortByGndx() EddyTry
{
std::vector<int> indicies(this->Size());
std::iota(indicies.begin(),indicies.end(),0);
std::sort(indicies.begin(),indicies.end(),
[this](int a, int b){ return(this->Gndx(a) < this->Gndx(b)); });
std::vector<std::shared_ptr<NEWIMAGE::volume<float> > > sptr_lst(this->Size());
std::vector<int> gndx_lst(this->Size());
for (unsigned int i=0; i<this->Size(); i++) {
sptr_lst[i] = this->_sptr_lst[indicies[i]];
gndx_lst[i] = this->_gndx_lst[indicies[i]];
}
this->_sptr_lst = sptr_lst;
this->_gndx_lst = gndx_lst;
} EddyCatch
fmriPredictor.h 0 → 100644
View file @ b9e18ed6
/*! \file fmriPredictor.h
\brief Contains declaration of class for making predictions about fmri data.
\author Jesper Andersson
\version 1.0b, Sep., 2012.
*/
//
// fmmriPredictor.h
//
// Jesper Andersson, FMRIB Image Analysis Group
//
// Copyright (C) 2022 University of Oxford
//
#ifndef fmriPredictor_h
#define fmriPredictor_h
#include <cstdlib>
#include <string>
#include <vector>
#include <cmath>
#include <memory>
#include "armawrap/newmat.h"
#include "newimage/newimageall.h"
#include "miscmaths/miscmaths.h"
#include "EddyHelperClasses.h"
#include "DWIPredictionMaker.h"
#include "KMatrix.h"
#include "HyParEstimator.h"
namespace EDDY {
/****************************************************************//**
*
* \brief
*
*
*
********************************************************************/
class fmriPredictor : public DWIPredictionMaker
{
public:
/// Default Constructor
fmriPredictor(const std::shared_ptr<const KMatrix>& Kmat,
const std::shared_ptr<const HyParEstimator>& hpe) EddyTry : _Kmats(1,Kmat->Clone()), _hpe(hpe->Clone()), _tr(-1.0) {} EddyCatch
~fmriPredictor() {}
/// Returns prediction for point given by indx.
virtual NEWIMAGE::volume<float> Predict(unsigned int indx, bool exclude=false) const;
/// Returns prediction for point given by indx.
virtual NEWIMAGE::volume<float> Predict(unsigned int indx, bool exclude=false);
/// Returns prediction for point given by indx. This is used only as a means of directly comparing CPU and GPU outputs.
virtual NEWIMAGE::volume<float> PredictCPU(unsigned int indx, bool exclude=false);
/// Returns predictions for points given by indicies
virtual std::vector<NEWIMAGE::volume<float> > Predict(const std::vector<unsigned int>& indicies, bool exclude=false);
/// Returns input data for point given by indx
virtual NEWIMAGE::volume<float> InputData(unsigned int indx) const;
/// Returns input data for points given by indicies
virtual std::vector<NEWIMAGE::volume<float> > InputData(const std::vector<unsigned int>& indicies) const;
/// Returns variance of prediction for point given by indx.
virtual double PredictionVariance(unsigned int indx, bool exclude=false);
/// Returns measurement-error variance for point given by indx.
virtual double ErrorVariance(unsigned int indx) const;
/// Returns true if all data has been loaded
virtual bool IsPopulated() const { return(is_populated()); }
/// Indicates if it is ready to make predictions.
virtual bool IsValid() const EddyTry { return(IsPopulated() && _lak && _Kmats[0]->IsValid()); } EddyCatch
/// Set total number of scans to be loaded
virtual void SetNoOfScans(unsigned int n);
/// Set a point given by indx. This function is thread safe as long as different threads set different points.
virtual void SetScan(const NEWIMAGE::volume<float>& scan, // _May_ be thread safe if used "sensibly"
const DiffPara& dp,
unsigned int indx,
unsigned int sess);
/// Set a point given by indx. This function is thread safe as long as different threads set different points.
virtual void SetScan(const NEWIMAGE::volume<float>& scan, // _May_ be thread safe if used "sensibly"
const DiffPara& dp,
unsigned int indx) EddyTry { SetScan(scan,dp,indx,0); } EddyCatch
/// Returns the number of hyperparameters for the model
virtual unsigned int NoOfHyperPar() const EddyTry { return(_Kmats[0]->NoOfHyperPar()); } EddyCatch
/// Returns the hyperparameters for the model
virtual std::vector<double> GetHyperPar() const EddyTry { return(_Kmats[0]->GetHyperPar()); } EddyCatch
/// Evaluates the model so as to make the predictor ready to make predictions.
virtual void EvaluateModel(const NEWIMAGE::volume<float>& mask, float fwhm, bool verbose=false);
/// Evaluates the model so as to make the predictor ready to make predictions.
virtual void EvaluateModel(const NEWIMAGE::volume<float>& mask, bool verbose=false) EddyTry { EvaluateModel(mask,0.0,verbose); } EddyCatch
/// Writes internal content to disk for debug purposes
virtual void WriteImageData(const std::string& fname) const;
virtual void WriteMetaData(const std::string& fname) const;
virtual void Write(const std::string& fname) const EddyTry { WriteImageData(fname); WriteMetaData(fname); } EddyCatch
private:
// Couple of helper classes
class sess_index {
public:
sess_index() : _sess(-1), _sindx(-1) {}
int _sess; // Session
int _sindx; // Index within session
};
class ptr_index_list {
public:
void PushBack(const NEWIMAGE::volume<float>& scan,
unsigned int indx) EddyTry {
_sptr_lst.push_back(std::make_shared<NEWIMAGE::volume<float> >(scan));
_gndx_lst.push_back(static_cast<int>(indx));
} EddyCatch
unsigned int Size() const { return(_sptr_lst.size()); }
std::shared_ptr<NEWIMAGE::volume<float> > SPtr(unsigned int i) const EddyTry {
if (i >= _sptr_lst.size()) throw EddyException("fmriPredictor::Sptr: Index out of bounds");
return(_sptr_lst[i]);
} EddyCatch
std::shared_ptr<NEWIMAGE::volume<float> >& SPtr(unsigned int i) EddyTry {
if (i >= _sptr_lst.size()) throw EddyException("fmriPredictor::Sptr: Index out of bounds");
return(_sptr_lst[i]);
} EddyCatch
int Gndx(unsigned int i) const EddyTry {
if (i >= _sptr_lst.size()) throw EddyException("fmriPredictor::Gndx: Index out of bounds");
return(_gndx_lst[i]);
} EddyCatch
const std::vector<std::shared_ptr<NEWIMAGE::volume<float> > >& SPtr_list() const EddyTry { return(_sptr_lst); } EddyCatch
void SortByGndx(); // Sort both _sptr_lst and _gndx_lst in ascending order of gndx
private:
std::vector<std::shared_ptr<NEWIMAGE::volume<float> > > _sptr_lst;
std::vector<int> _gndx_lst;
};
// Member variables
std::vector<sess_index> _glist; /// List indexed with "global" index
std::vector<ptr_index_list> _slist; /// List indexed with session index
std::vector<std::shared_ptr<KMatrix> > _Kmats; /// K-matrices
std::shared_ptr<HyParEstimator> _hpe;
std::vector<std::shared_ptr<NEWIMAGE::volume<float> > > _mptrs; /// Pointers to mean images
double _tr; /// Repetition time
bool _lak; /// Lists Are Kosher?
std::mutex _set_mut; /// Mutex for SetNoOfScans and SetScan
// Private utility functions
void mean_correct();
bool is_populated() const;
/// Return session index for image with global index 'gindx'
int session(int gindx) const EddyTry {
if (!this->is_populated()) throw EddyException("fmriPredictor::no_of_scans_in_same_session: Predictor not yet populated");
if (gindx >= _glist.size()) throw EddyException("fmriPredictor::no_of_scans_in_same_session: Index out of bounds");
return(_glist[gindx]._sess);
} EddyCatch
/// Return true of all the global indicies in 'indicies' belong to the same session
bool all_in_same_session(const std::vector<unsigned int>& indicies) const EddyTry {
int sess = this->session(indicies[0]);
for (unsigned int i=1; i<indicies.size(); i++) if (sess != this->session(indicies[i])) return(false);
return(true);
} EddyCatch
/// Returns a ptr to first image. Useful for getting size and properties.
std::shared_ptr<const NEWIMAGE::volume<float> > first_imptr() const EddyTry {
if (!this->is_populated()) throw EddyException("fmriPredictor::first_imptr: Predictor not yet populated");
return(_slist[0].SPtr(0));
} EddyCatch
/// No of scans in same session as that given by global index gindx
unsigned int no_of_scans_in_same_session(int gindx) const EddyTry {
if (!this->is_populated()) throw EddyException("fmriPredictor::no_of_scans_in_same_session: Predictor not yet populated");
if (gindx >= _glist.size()) throw EddyException("fmriPredictor::no_of_scans_in_same_session: Index out of bounds");
return(_slist[_glist[gindx]._sess].Size());
} EddyCatch
/// Returns index within session for scan with global index "gindx"
int index_in_session(int gindx) const EddyTry {
if (!this->is_populated()) throw EddyException("fmriPredictor::index_in_session: Predictor not yet populated");
if (gindx >= _glist.size()) throw EddyException("fmriPredictor::index_in_session: Index out of bounds");
return(_glist[gindx]._sindx);
} EddyCatch
/// Returns image ptr #i from the same session as gindx
std::shared_ptr<const NEWIMAGE::volume<float> > imptr_from_same_session(int gindx, int i) const EddyTry {
if (!this->is_populated()) throw EddyException("fmriPredictor::gindx_and_imptr_from_same_session: Predictor not yet populated");
if (gindx >= _glist.size()) throw EddyException("fmriPredictor::gindx_and_imptr_from_same_session: gindx out of bounds");
if (i >= _slist[_glist[gindx]._sess].Size()) throw EddyException("fmriPredictor::gindx_and_imptr_from_same_session: i out of bounds");
return(_slist[_glist[gindx]._sess].SPtr(i));
} EddyCatch
/// Returns mean_image ptr from the same session as gindx
std::shared_ptr<const NEWIMAGE::volume<float> > meanptr_from_same_session(int gindx) const EddyTry {
if (!this->is_populated()) throw EddyException("fmriPredictor::gindx_and_imptr_from_same_session: Predictor not yet populated");
if (gindx >= _glist.size()) throw EddyException("fmriPredictor::gindx_and_imptr_from_same_session: gindx out of bounds");
return(_mptrs[_glist[gindx]._sess]);
} EddyCatch
/*
/// Returns all global indicies (including gindx) from the same session as gindx
std::vector<int> global_indicies_in_same_session(int gindx) const EddyTry {
if (!this->is_populated()) throw EddyException("fmriPredictor::global_indicies_in_same_session: Predictor not yet populated");
if (gindx >= _glist.size()) throw EddyException("fmriPredictor::global_indicies_in_same_session: Index out of bounds");
std::vector<int> rval(_slist[_glist[gindx]._sess].Size());
for (int i=0; i<rval.size(); i++) rval[i] = _slist[_glist[gindx]._sess].Gndx(i);
return(rval);
} EddyCatch
/// Returns global index and image ptr #i from the same session as gindx
std::tuple<int, std::shared_ptr<NEWIMAGE::volume<float> > > gindx_and_imptr_from_same_session(int gindx, int i) EddyTry {
if (!this->is_populated()) throw EddyException("fmriPredictor::gindx_and_imptr_from_same_session: Predictor not yet populated");
if (gindx >= _glist.size()) throw EddyException("fmriPredictor::gindx_and_imptr_from_same_session: gindx out of bounds");
if (i >= _slist[_glist[gindx]._sess].Size()) throw EddyException("fmriPredictor::gindx_and_imptr_from_same_session: i out of bounds");
return(std::make_tuple(_slist[_glist[gindx]._sess].Gndx(i),_slist[_glist[gindx]._sess].Sptr(i)));
}
*/
/// Total number of scans
unsigned int no_of_scans() const EddyTry {
if (!this->is_populated()) throw EddyException("fmriPredictor::no_of_scans: Predictor not yet populated");
return(_glist.size());
} EddyCatch
/// Total number of scans in session
unsigned int no_of_scans(unsigned int session) const EddyTry {
if (!this->is_populated()) throw EddyException("fmriPredictor::no_of_scans: Predictor not yet populated");
if (session >= _slist.size()) throw EddyException("fmriPredictor::no_of_scans: Index out of bounds");
return(_slist[session].Size());
} EddyCatch
bool same_no_of_scans_in_all_sessions() const;
bool lists_are_kosher();
bool glist_has_duplicates() const;
void predict_image_cpu(// Input
unsigned int indx,
bool exclude,
const arma::rowvec& pv,
// Output
NEWIMAGE::volume<float>& pi) const;
void predict_images_cpu(// Input
const std::vector<unsigned int>& indicies,
bool exclude,
const std::vector<arma::rowvec>& pvecs,
// Output
std::vector<NEWIMAGE::volume<float> >& pi) const;
#ifdef COMPILE_GPU
void predict_image_gpu(unsigned int indx,
bool excl,
const arma::rowvec& pv,
NEWIMAGE::volume<float>& ima) const;
void predict_images_gpu(// Input
const std::vector<unsigned int>& indicies,
bool exclude,
const std::vector<arma::rowvec>& pvecs,
// Output
std::vector<NEWIMAGE::volume<float> >& pi) const;
#endif
bool get_y(// Input
unsigned int i, unsigned int j, unsigned int k, unsigned int indx, bool exclude,
// Output
arma::colvec& y) const;
};
} // End namespace EDDY
#endif // End #ifndef fmriPredictor_h
log.log 0 → 100644
View file @ b9e18ed6
This source diff could not be displayed because it is too large. You can view the blob instead.
metadata_conda_debug.yaml 0 → 100644
View file @ b9e18ed6
package:
name: fsl-eddy
version: '2401.2'
source:
git_url: https://git.fmrib.ox.ac.uk/fsl/eddy.git
git_rev: '2401.2'
build:
number: '5'
run_exports:
strong:
- fsl-eddy
requirements:
host:
- libboost-headers
- fsl-base
- fsl-newnifti 4
- fsl-armawrap 0
- fsl-basisfield 2203
- fsl-cprob 2111
- fsl-miscmaths 2412
- fsl-newimage 2501
- fsl-topup 2203
- fsl-utils 2412
- fsl-warpfns 2501
- fsl-znzlib 2111
- nlohmann_json 3
- libblas
- liblapack
- zlib
build:
- gxx_linux-64
- make
- sysroot_linux-64 2.17.*
extra:
final: false
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