double evolve_D(struct network * network, long i, long j, long k, const double * y) { long compartment_offset = network->neurons[i]->compartments[j]->ode_system_offset; long link_offset = network->neurons[i]->compartments[j]->links[k]->ode_system_offset; struct neuron_params * params = network->neurons[i]->params; double tau_d = params->values[37]; double B = y[link_offset + 3]; double D = y[link_offset + 4]; return (d_sigma(network,B,i,j,k) - D)/tau_d; }
__host__ void calc_chisq_wrapper_c(const double *ph_model, const double *ph_obs, const double *ph_sigma, const int num_data, const int num_model, double *ph_chisq ) { int gpuid = ebf::init_cuda(); // put vectors in thrust format from raw points int num = num_data*num_model; thrust::host_vector<double> h_model(ph_model,ph_model+num); thrust::host_vector<double> h_obs(ph_obs,ph_obs+num_data); thrust::host_vector<double> h_sigma(ph_sigma,ph_sigma+num_data); thrust::host_vector<double> h_chisq(ph_chisq,ph_chisq+num_model); thrust::counting_iterator<int> index_begin(0); thrust::counting_iterator<int> index_end(num); if(gpuid>=0) { // allocate mem on GPU thrust::device_vector<double> d_model(num); thrust::device_vector<double> d_obs(num_data); thrust::device_vector<double> d_sigma(num_data); thrust::device_vector<double> d_chisq(num_model); // transfer input params to GPU // start_timer_upload(); d_model = h_model; d_obs = h_obs; d_sigma = h_sigma; // stop_timer_upload(); // distribute the computation to the GPU ebf::sync(); // start_timer_kernel(); typedef thrust::counting_iterator<int> CountIntIter; typedef thrust::transform_iterator<modulo_stride_functor,CountIntIter> ObsIdxIter; typedef thrust::permutation_iterator<thrust::device_vector<double>::iterator, ObsIdxIter> ObsIter; typedef thrust::tuple<thrust::device_vector<double>::iterator, ObsIter, ObsIter> ModelObsSigmaIteratorTuple; typedef thrust::zip_iterator<ModelObsSigmaIteratorTuple> ModelObsSigmaZipIter; ModelObsSigmaZipIter input_begin = thrust::make_zip_iterator(thrust::make_tuple( d_model.begin(), thrust::make_permutation_iterator(d_obs.begin(),thrust::make_transform_iterator(index_begin,modulo_stride_functor(num_data))), thrust::make_permutation_iterator(d_sigma.begin(),thrust::make_transform_iterator(index_begin,modulo_stride_functor(num_data))) )); for(int m=0;m<num_model;++m) { d_chisq[m] = thrust::transform_reduce( input_begin+m*num_data, input_begin+((m+1)*num_data), calc_ressq_functor(), 0., thrust::plus<double>() ); } ebf::sync(); // stop_timer_kernel(); // transfer results back to host // start_timer_download(); thrust::copy(d_chisq.begin(),d_chisq.end(),ph_chisq); // stop_timer_download(); } else { // distribute the computation to the CPU typedef thrust::counting_iterator<int> CountIntIter; typedef thrust::transform_iterator<modulo_stride_functor,CountIntIter> ObsIdxIter; typedef thrust::permutation_iterator<thrust::host_vector<double>::iterator, ObsIdxIter> ObsIter; typedef thrust::tuple<thrust::host_vector<double>::iterator, ObsIter, ObsIter> ModelObsSigmaIteratorTuple; typedef thrust::zip_iterator<ModelObsSigmaIteratorTuple> ModelObsSigmaZipIter; ModelObsSigmaZipIter input_begin = thrust::make_zip_iterator(thrust::make_tuple( h_model.begin(), thrust::make_permutation_iterator(h_obs.begin(),thrust::make_transform_iterator(index_begin,modulo_stride_functor(num_data))), thrust::make_permutation_iterator(h_sigma.begin(),thrust::make_transform_iterator(index_begin,modulo_stride_functor(num_data))) )); // start_timer_kernel(); for(int m=0;m<num_model;++m) { h_chisq[m] = thrust::transform_reduce( input_begin+m*num_data, input_begin+((m+1)*num_data), calc_ressq_functor(), 0., thrust::plus<double>() ); } // stop_timer_kernel(); } }