diff --git a/bioem.cpp b/bioem.cpp index 55bc03788d9b591ef9ec3d743961ac8d2a843274..7c4941c7328771f8f2135167711282ccb957bc81 100644 --- a/bioem.cpp +++ b/bioem.cpp @@ -270,7 +270,7 @@ int bioem::configure(int ac, char* av[]) printf("Time Precalculate %f\n", timer.GetCurrentElapsedTime()); timer.ResetStart(); } - pProb.init(RefMap.ntotRefMap, param.nTotGridAngles, *this); + pProb.init(RefMap.ntotRefMap, param.nTotGridAngles, param.nTotCC, *this); if (DebugOutput >= 2 && mpi_rank == 0) { diff --git a/bioem_cuda.cu b/bioem_cuda.cu index 087dfafc0719ccaff3553e57e2a3f3f643d20a51..891a28e916e9fcb937b3b0790831e5df50b9d4ab 100644 --- a/bioem_cuda.cu +++ b/bioem_cuda.cu @@ -397,8 +397,7 @@ int bioem_cuda::deviceInit() gpumap->sum_RefMap = sum; gpumap->sumsquare_RefMap = sumsquare; - checkCudaErrors(cudaMalloc(&pProb_memory, pProb_device.get_size(RefMap.ntotRefMap, param.nTotGridAngles, param.param_device.writeAngles))); - checkCudaErrors(cudaMalloc(&pProb_memory, pProb_device.get_sizeCC(RefMap.ntotRefMap, param.nTotCC, param.param_device.writeCC))); + checkCudaErrors(cudaMalloc(&pProb_memory, pProb_device.get_size(RefMap.ntotRefMap, param.nTotGridAngles, param.nTotCC, param.param_device.writeAngles, param.param_device.writeCC))); for (int i = 0; i < 2; i++) { @@ -484,15 +483,14 @@ int bioem_cuda::deviceStartRun() { maxRef = (size_t) RefMap.ntotRefMap * (size_t) GPUWorkload / 100; pProb_host = new bioem_Probability; - pProb_host->init(maxRef, param.nTotGridAngles, *this); - pProb_host->initCC(maxRef, param.nTotCC, *this); + pProb_host->init(maxRef, param.nTotGridAngles, param.nTotCC, *this); pProb_host->copyFrom(&pProb, *this); } pProb_device = *pProb_host; pProb_device.ptr = pProb_memory; pProb_device.set_pointers(); - checkCudaErrors(cudaMemcpyAsync(pProb_device.ptr, pProb_host->ptr, pProb_host->get_size(maxRef, param.nTotGridAngles, param.param_device.writeAngles), cudaMemcpyHostToDevice, cudaStream[0])); + checkCudaErrors(cudaMemcpyAsync(pProb_device.ptr, pProb_host->ptr, pProb_host->get_size(maxRef, param.nTotGridAngles, param.nTotCC, param.param_device.writeAngles, param.param_device.writeCC), cudaMemcpyHostToDevice, cudaStream[0])); if (FFTAlgo) { @@ -527,7 +525,7 @@ int bioem_cuda::deviceStartRun() int bioem_cuda::deviceFinishRun() { if (GPUAsync) cudaStreamSynchronize(cudaStream[0]); - checkCudaErrors(cudaMemcpyAsync(pProb_host->ptr, pProb_device.ptr, pProb_host->get_size(maxRef, param.nTotGridAngles, param.param_device.writeAngles), cudaMemcpyDeviceToHost, cudaStream[0])); + checkCudaErrors(cudaMemcpyAsync(pProb_host->ptr, pProb_device.ptr, pProb_host->get_size(maxRef, param.nTotGridAngles, param.nTotCC, param.param_device.writeAngles, param.param_device.writeCC), cudaMemcpyDeviceToHost, cudaStream[0])); if (FFTAlgo) { diff --git a/include/map.h b/include/map.h index 3b67da59e1c0e59cdad42a58c94a728d9b2af0f1..5500f239286f541c88d5471d77331d72113cabeb 100644 --- a/include/map.h +++ b/include/map.h @@ -91,7 +91,7 @@ public: { public: int max_prob_cent_x, max_prob_cent_y, max_prob_orient, max_prob_conv; - myfloat_t max_prob_norm,max_prob_mu; + myfloat_t max_prob_norm,max_prob_mu; } max; }; @@ -105,7 +105,7 @@ public: class bioem_Probability_cc { public: - myfloat_t forCC; + myfloat_t forCC; }; class bioem_Probability @@ -113,39 +113,33 @@ class bioem_Probability public: int nMaps; int nAngles; - int nCC; + int nCC; __device__ __host__ bioem_Probability_map& getProbMap(int map) {return(ptr_map[map]);} __device__ __host__ bioem_Probability_angle& getProbAngle(int map, int angle) {return(ptr_angle[angle * nMaps + map]);} - __device__ __host__ bioem_Probability_cc& getProbCC(int map, int cc) {return(ptr_cc[cc * nMaps + map]);} + __device__ __host__ bioem_Probability_cc& getProbCC(int map, int cc) {return(ptr_cc[cc * nMaps + map]);} void* ptr; bioem_Probability_map* ptr_map; bioem_Probability_angle* ptr_angle; - bioem_Probability_cc* ptr_cc; + bioem_Probability_cc* ptr_cc; - static size_t get_size(size_t maps, size_t angles, bool writeAngles) + static size_t get_size(size_t maps, size_t angles, size_t cc, bool writeAngles, bool writeCC) { size_t size = sizeof(bioem_Probability_map); if (writeAngles) size += angles * sizeof(bioem_Probability_angle); + if (writeCC) size += cc * sizeof(bioem_Probability_cc); return(maps * size); } - static size_t get_sizeCC(size_t maps, size_t cc, bool writeCC) - { - size_t size = sizeof(bioem_Probability_map); - if (writeCC) size += cc * sizeof(bioem_Probability_cc) ; - return(maps * size); - } - - void init(size_t maps, size_t angles, bioem& bio); - void initCC(size_t maps, size_t cc, bioem& bio); + + void init(size_t maps, size_t angles, size_t cc, bioem& bio); void copyFrom(bioem_Probability* from, bioem& bio); void set_pointers() { ptr_map = (bioem_Probability_map*) ptr; ptr_angle = (bioem_Probability_angle*) (&ptr_map[nMaps]); - ptr_cc = (bioem_Probability_cc*) (&ptr_map[nMaps]); + ptr_cc = (bioem_Probability_cc*) (&ptr_angle[nMaps * nAngles]); } }; diff --git a/map.cpp b/map.cpp index 3c5aaf86e2bb740dbfc66f8558dd24895597a729..dedb1ac94c3f1e0c5ef3183281dbe1d307fcd7b1 100644 --- a/map.cpp +++ b/map.cpp @@ -280,23 +280,15 @@ int bioem_RefMap::precalculate(bioem_param& param, bioem& bio) return(0); } -void bioem_Probability::init(size_t maps, size_t angles, bioem& bio) +void bioem_Probability::init(size_t maps, size_t angles, size_t cc, bioem& bio) { nMaps = maps; nAngles = angles; - ptr = bio.malloc_device_host(get_size(maps, angles, bio.param.param_device.writeAngles)); + nCC = cc; + ptr = bio.malloc_device_host(get_size(maps, angles, cc, bio.param.param_device.writeAngles, bio.param.param_device.writeCC)); set_pointers(); } -void bioem_Probability::initCC(size_t maps, size_t cc, bioem& bio) -{ - nMaps = maps; - nCC = cc; - ptr = bio.malloc_device_host(get_sizeCC(maps, cc, bio.param.param_device.writeCC)); - set_pointers(); -} - - void bioem_Probability::copyFrom(bioem_Probability* from, bioem& bio) { bioem_Probability_map& pProbMap = getProbMap(0); @@ -305,13 +297,23 @@ void bioem_Probability::copyFrom(bioem_Probability* from, bioem& bio) if (bio.param.param_device.writeAngles) { - for (int iOrient = 0; iOrient < bio.param.nTotGridAngles; iOrient ++) + for (int iOrient = 0; iOrient < nAngles; iOrient ++) { bioem_Probability_angle& pProbAngle = getProbAngle(0, iOrient); bioem_Probability_angle& pProbAngleFrom = from->getProbAngle(0, iOrient); memcpy(&pProbAngle, &pProbAngleFrom, from->nMaps * sizeof(bioem_Probability_angle)); } } + + if (bio.param.param_device.writeCC) + { + for (int iCC = 0; iCC < nCC; iCC ++) + { + bioem_Probability_cc& pProbCC = getProbCC(0, iCC); + bioem_Probability_cc& pProbCCFrom = from->getProbCC(0, iCC); + memcpy(&pProbCC, &pProbCCFrom, from->nMaps * sizeof(bioem_Probability_cc)); + } + } } int bioem_RefMap::read_MRC(const char* filename,bioem_param& param)