From 87b058d44107096d3c285a065ede2b94b587455a Mon Sep 17 00:00:00 2001 From: Moises Fernandez <moisesf@fmrib.ox.ac.uk> Date: Fri, 19 Apr 2013 11:13:07 +0000 Subject: [PATCH] Delete unnecessary params in Samples Record-Save --- CUDA/runmcmc.cu | 10 +++------- CUDA/runmcmc.h | 4 +--- CUDA/samples.cu | 20 ++++---------------- CUDA/samples.h | 3 +-- CUDA/xfibres_gpu.cuh | 7 +------ 5 files changed, 10 insertions(+), 34 deletions(-) diff --git a/CUDA/runmcmc.cu b/CUDA/runmcmc.cu index c09aa06..3fd3360 100644 --- a/CUDA/runmcmc.cu +++ b/CUDA/runmcmc.cu @@ -288,7 +288,6 @@ void runmcmc_record( //INPUT double seed, string output_file, //OUTPUT - thrust::device_vector<int>& multirecords_gpu, thrust::device_vector<float>& rf0_gpu, thrust::device_vector<float>& rtau_gpu, thrust::device_vector<float>& rs0_gpu, @@ -296,8 +295,7 @@ void runmcmc_record( //INPUT thrust::device_vector<float>& rdstd_gpu, thrust::device_vector<float>& rth_gpu, thrust::device_vector<float>& rph_gpu, - thrust::device_vector<float>& rf_gpu, - thrust::device_vector<float>& rlikelihood_energy_gpu) + thrust::device_vector<float>& rf_gpu) { xfibresOptions& opts = xfibresOptions::getInstance(); @@ -397,7 +395,6 @@ void runmcmc_record( //INPUT double *signals_ptr = thrust::raw_pointer_cast(signals_gpu.data()); double *isosignals_ptr = thrust::raw_pointer_cast(isosignals_gpu.data()); - int *multirecords_ptr = thrust::raw_pointer_cast(multirecords_gpu.data()); float *rf0_ptr = thrust::raw_pointer_cast(rf0_gpu.data()); float *rtau_ptr = thrust::raw_pointer_cast(rtau_gpu.data()); float *rs0_ptr = thrust::raw_pointer_cast(rs0_gpu.data()); @@ -406,7 +403,6 @@ void runmcmc_record( //INPUT float *rth_ptr = thrust::raw_pointer_cast(rth_gpu.data()); float *rph_ptr = thrust::raw_pointer_cast(rph_gpu.data()); float *rf_ptr = thrust::raw_pointer_cast(rf_gpu.data()); - float *rlikelihood_energy_ptr = thrust::raw_pointer_cast(rlikelihood_energy_gpu.data()); int amount_shared = (THREADS_BLOCK_MCMC+1)*sizeof(double) + (10*nfib + 2*nparams + 24)*sizeof(float) + (7*nfib + 18)*sizeof(int); @@ -434,7 +430,7 @@ void runmcmc_record( //INPUT gettimeofday(&t1,NULL); - runmcmc_record_kernel<<< Dim_Grid, Dim_Block, amount_shared >>>(datam_ptr, bvals_ptr, alpha_ptr, beta_ptr, fibres_ptr, multifibres_ptr, signals_ptr, isosignals_ptr, randomsN_ptr, randomsU_ptr, ndirections, nfib, nparams, opts.modelnum.value(), opts.fudge.value(), opts.f0.value(), opts.ardf0.value(), !opts.no_ard.value(), opts.rician.value(), opts.updateproposalevery.value(), iters_step, (i*iters_step), opts.nburn.value(), opts.sampleevery.value(), totalrecords, multirecords_ptr, rf0_ptr, rtau_ptr, rs0_ptr, rd_ptr, rdstd_ptr, rth_ptr, rph_ptr, rf_ptr, rlikelihood_energy_ptr); + runmcmc_record_kernel<<< Dim_Grid, Dim_Block, amount_shared >>>(datam_ptr, bvals_ptr, alpha_ptr, beta_ptr, fibres_ptr, multifibres_ptr, signals_ptr, isosignals_ptr, randomsN_ptr, randomsU_ptr, ndirections, nfib, nparams, opts.modelnum.value(), opts.fudge.value(), opts.f0.value(), opts.ardf0.value(), !opts.no_ard.value(), opts.rician.value(), opts.updateproposalevery.value(), iters_step, (i*iters_step), opts.nburn.value(), opts.sampleevery.value(), totalrecords, rf0_ptr, rtau_ptr, rs0_ptr, rd_ptr, rdstd_ptr, rth_ptr, rph_ptr, rf_ptr); sync_check("runmcmc_record_kernel"); gettimeofday(&t2,NULL); @@ -464,7 +460,7 @@ void runmcmc_record( //INPUT gettimeofday(&t1,NULL); if(nvox!=0){ - runmcmc_record_kernel<<< Dim_Grid, Dim_Block, amount_shared >>>(datam_ptr, bvals_ptr, alpha_ptr, beta_ptr, fibres_ptr, multifibres_ptr, signals_ptr, isosignals_ptr, randomsN_ptr, randomsU_ptr, ndirections, nfib, nparams, opts.modelnum.value(), opts.fudge.value(), opts.f0.value(), opts.ardf0.value(), !opts.no_ard.value(), opts.rician.value(), opts.updateproposalevery.value(), last_step, (steps*iters_step), opts.nburn.value(), opts.sampleevery.value(), totalrecords, multirecords_ptr, rf0_ptr, rtau_ptr, rs0_ptr, rd_ptr, rdstd_ptr, rth_ptr, rph_ptr, rf_ptr, rlikelihood_energy_ptr); + runmcmc_record_kernel<<< Dim_Grid, Dim_Block, amount_shared >>>(datam_ptr, bvals_ptr, alpha_ptr, beta_ptr, fibres_ptr, multifibres_ptr, signals_ptr, isosignals_ptr, randomsN_ptr, randomsU_ptr, ndirections, nfib, nparams, opts.modelnum.value(), opts.fudge.value(), opts.f0.value(), opts.ardf0.value(), !opts.no_ard.value(), opts.rician.value(), opts.updateproposalevery.value(), last_step, (steps*iters_step), opts.nburn.value(), opts.sampleevery.value(), totalrecords,rf0_ptr, rtau_ptr, rs0_ptr, rd_ptr, rdstd_ptr, rth_ptr, rph_ptr, rf_ptr); sync_check("runmcmc_record_kernel"); } diff --git a/CUDA/runmcmc.h b/CUDA/runmcmc.h index c918496..5d87331 100644 --- a/CUDA/runmcmc.h +++ b/CUDA/runmcmc.h @@ -53,7 +53,6 @@ void runmcmc_record( //INPUT double seed, string output_file, //OUTPUT - thrust::device_vector<int>& multirecords_gpu, thrust::device_vector<float>& rf0_gpu, thrust::device_vector<float>& rtau_gpu, thrust::device_vector<float>& rs0_gpu, @@ -61,5 +60,4 @@ void runmcmc_record( //INPUT thrust::device_vector<float>& rdstd_gpu, thrust::device_vector<float>& rth_gpu, thrust::device_vector<float>& rph_gpu, - thrust::device_vector<float>& rf_gpu, - thrust::device_vector<float>& rlikelihood_energy_gpu); + thrust::device_vector<float>& rf_gpu); diff --git a/CUDA/samples.cu b/CUDA/samples.cu index 11cb4ee..f4f4886 100644 --- a/CUDA/samples.cu +++ b/CUDA/samples.cu @@ -17,7 +17,7 @@ using namespace Xfibres; // MCMC SAMPLE STORAGE //////////////////////////////////////////// -Samples::Samples(int nvoxels,int nmeasures): +Samples::Samples(int nvoxels,int nsamples): opts(xfibresOptions::getInstance()){ /////////////// GPU version ///////////////////// @@ -33,23 +33,11 @@ opts(xfibresOptions::getInstance()){ m_sum_lam=new vector<float> [nvoxels]; //////////////////////////////////////////////// - //m_beenhere=m_vol2matrixkey*0; - int count=0; - int nsamples=0; - - for(int i=0;i<opts.njumps.value();i++){ - count++; - if(count==opts.sampleevery.value()){ - count=0;nsamples++; - } - } - m_dsamples.ReSize(nsamples,nvoxels); m_dsamples=0; m_S0samples.ReSize(nsamples,nvoxels); m_S0samples=0; - m_lik_energy.ReSize(nsamples,nvoxels); - + m_mean_dsamples.ReSize(nvoxels); m_mean_dsamples=0; m_mean_S0samples.ReSize(nvoxels); @@ -133,7 +121,7 @@ opts(xfibresOptions::getInstance()){ } //new version for GPU -void Samples::record(float rd,float rf0,float rtau,float rdstd,float rs0,float rlikelihood_energy, float *rth,float *rph, float *rf, int vox, int samp){ +void Samples::record(float rd,float rf0,float rtau,float rdstd,float rs0,float *rth,float *rph, float *rf, int vox, int samp){ m_dsamples(samp,vox)=rd; m_sum_d[vox-1]+=rd; @@ -151,7 +139,7 @@ void Samples::record(float rd,float rf0,float rtau,float rdstd,float rs0,float r m_S0samples(samp,vox)=rs0; m_sum_S0[vox-1]+=rs0; - m_lik_energy(samp,vox)=rlikelihood_energy; + for(int f=0;f<opts.nfibres.value();f++){ float th=rth[f]; float ph=rph[f]; diff --git a/CUDA/samples.h b/CUDA/samples.h index 03ee1b1..abb675d 100644 --- a/CUDA/samples.h +++ b/CUDA/samples.h @@ -23,7 +23,6 @@ class Samples{ Matrix m_d_stdsamples; Matrix m_S0samples; Matrix m_f0samples; - Matrix m_lik_energy; // // storing signal // Matrix m_mean_sig; @@ -74,7 +73,7 @@ class Samples{ Samples(int nvoxels,int nmeasures); - void record(float rd,float rf0,float rtau,float rdstd,float rs0,float rlikelihood_energy, float *rth,float *rph, float *rf, int vox, int samp); + void record(float rd,float rf0,float rtau,float rdstd,float rs0,float *rth,float *rph, float *rf, int vox, int samp); void finish_voxel(int vox); diff --git a/CUDA/xfibres_gpu.cuh b/CUDA/xfibres_gpu.cuh index 07f6d2c..eb88dd1 100644 --- a/CUDA/xfibres_gpu.cuh +++ b/CUDA/xfibres_gpu.cuh @@ -91,7 +91,6 @@ void prepare_data_gpu_MCMC( //INPUT void prepare_data_gpu_MCMC_record( //INPUT int nvox, //OUTPUT - thrust::device_vector<int>& multirecords_gpu, thrust::device_vector<float>& rf0_gpu, thrust::device_vector<float>& rtau_gpu, thrust::device_vector<float>& rs0_gpu, @@ -99,12 +98,10 @@ void prepare_data_gpu_MCMC_record( //INPUT thrust::device_vector<float>& rdstd_gpu, thrust::device_vector<float>& rth_gpu, thrust::device_vector<float>& rph_gpu, - thrust::device_vector<float>& rf_gpu, - thrust::device_vector<float>& rlikelihood_energy_gpu); + thrust::device_vector<float>& rf_gpu); //implemented and used in xfibres_gpu.cu void record_finish_voxels( //INPUT - thrust::device_vector<int>& multirecords_gpu, thrust::device_vector<float>& rf0_gpu, thrust::device_vector<float>& rtau_gpu, thrust::device_vector<float>& rs0_gpu, @@ -113,8 +110,6 @@ void record_finish_voxels( //INPUT thrust::device_vector<float>& rth_gpu, thrust::device_vector<float>& rph_gpu, thrust::device_vector<float>& rf_gpu, - thrust::device_vector<float>& rlikelihood_energy_gpu, int nvox, - int ndirections, int idpart); -- GitLab