Skip to content

Commit c20fac6

Browse files
krzysgkrzysg
committed
localScaleTemp is not anymore passed to GpuProcessinTask, CPU memory init in APRConverter for GPU not needed anymore
1 parent 3a11834 commit c20fac6

7 files changed

Lines changed: 44 additions & 55 deletions

‎src/algorithm/APRConverter.hpp‎

Lines changed: 2 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -152,7 +152,7 @@ void APRConverter<ImageType>::initPipelineMemory(int y_num,int x_num,int z_num){
152152

153153
float not_needed;
154154
std::vector<int> var_win;
155-
iLocalIntensityScale.get_window_alt(not_needed, var_win, par, grad_temp);
155+
iLocalIntensityScale.get_window_alt(not_needed, var_win, par, grad_temp.getDimension());
156156

157157
int padding_y = 2*std::max(var_win[0],var_win[3]);
158158
int padding_x = 2*std::max(var_win[1],var_win[4]);
@@ -434,7 +434,6 @@ inline bool APRConverter<ImageType>::get_apr_cuda_multistreams(std::vector<APR*>
434434
for (auto apr : aAPRs) {
435435
if (!initPipelineAPR(*apr, *input_image)) return false;
436436
}
437-
initPipelineMemory(input_image->y_num, input_image->x_num, input_image->z_num);
438437

439438
// Create a temporary image for each stream
440439
std::vector<PixelData<ImageType>> tempImages;
@@ -453,7 +452,7 @@ inline bool APRConverter<ImageType>::get_apr_cuda_multistreams(std::vector<APR*>
453452
t.start_timer("Creating GPTS");
454453
std::vector<std::future<void>> gpts_futures; gpts_futures.resize(numOfStreams);
455454
for (int i = 0; i < numOfStreams; ++i) {
456-
gpts.emplace_back(GpuProcessingTask<ImageType>(tempImages[i], local_scale_temp, par, aAPRs[0]->level_max()));
455+
gpts.emplace_back(GpuProcessingTask<ImageType>(tempImages[i], par, aAPRs[0]->level_max()));
457456
}
458457
t.stop_timer();
459458

‎src/algorithm/ComputeGradientCuda.cu‎

Lines changed: 25 additions & 32 deletions
Original file line numberDiff line numberDiff line change
@@ -172,7 +172,7 @@ namespace {
172172
}
173173

174174
template <typename ImgType>
175-
void getGradientCuda(const PixelData<ImgType> &image, PixelData<float> &local_scale_temp,
175+
void getGradientCuda(const PixelData<ImgType> &image,
176176
ImgType *cudaImage, ImgType *cudaGrad, float *cudalocal_scale_temp,
177177
BsplineParamsCuda &px, BsplineParamsCuda &py, BsplineParamsCuda &pz, float *boundary,
178178
bool &isErrorDetected, ScopedCudaMemHandler<bool *, JUST_ALLOC>& isErrorDetectedCuda,
@@ -196,13 +196,14 @@ void getGradientCuda(const PixelData<ImgType> &image, PixelData<float> &local_sc
196196
"try squashing the input image to a narrower range or use APRConverter<float>");
197197
}
198198
}
199-
runKernelGradient(cudaImage, cudaGrad, image.getDimension(), local_scale_temp.getDimension(), par.dx, par.dy, par.dz, aStream);
199+
auto localScaleTempDim = image.getDimensionDS(); // size of downsampled input image
200+
runKernelGradient(cudaImage, cudaGrad, image.getDimension(), localScaleTempDim, par.dx, par.dy, par.dz, aStream);
200201
runDownsampleMean(cudaImage, cudalocal_scale_temp, image.x_num, image.y_num, image.z_num, aStream);
201202

202203
if (par.lambda > 0) {
203-
if (image.y_num > 2) runInvBsplineYdir(cudalocal_scale_temp, local_scale_temp.x_num, local_scale_temp.y_num, local_scale_temp.z_num, aStream);
204-
if (image.x_num > 2) runInvBsplineXdir(cudalocal_scale_temp, local_scale_temp.x_num, local_scale_temp.y_num, local_scale_temp.z_num, aStream);
205-
if (image.z_num > 2) runInvBsplineZdir(cudalocal_scale_temp, local_scale_temp.x_num, local_scale_temp.y_num, local_scale_temp.z_num, aStream);
204+
if (image.y_num > 2) runInvBsplineYdir(cudalocal_scale_temp, localScaleTempDim.x, localScaleTempDim.y, localScaleTempDim.z, aStream);
205+
if (image.x_num > 2) runInvBsplineXdir(cudalocal_scale_temp, localScaleTempDim.x, localScaleTempDim.y, localScaleTempDim.z, aStream);
206+
if (image.z_num > 2) runInvBsplineZdir(cudalocal_scale_temp, localScaleTempDim.x, localScaleTempDim.y, localScaleTempDim.z, aStream);
206207
}
207208
}
208209

@@ -429,7 +430,7 @@ class GpuProcessingTask<U>::GpuProcessingTaskImpl {
429430

430431
// input data
431432
const PixelData<ImgType> &iCpuImage;
432-
PixelData<float> &iCpuLevels;
433+
// PixelData<float> &iCpuLevels;
433434
const APRParameters &iParameters;
434435
GenInfo iAprInfo;
435436
float iBsplineOffset = 0;
@@ -438,9 +439,9 @@ class GpuProcessingTask<U>::GpuProcessingTaskImpl {
438439
// cuda stuff - memory and stream to be used
439440
ScopedCudaMemHandler<const PixelData<ImgType>, JUST_ALLOC> image;
440441
ScopedCudaMemHandler<const PixelData<ImgType>, JUST_ALLOC> imageSampling;
441-
ScopedCudaMemHandler<PixelData<ImgType>, JUST_ALLOC> gradient;
442-
ScopedCudaMemHandler<PixelData<float>, JUST_ALLOC> local_scale_temp;
443-
ScopedCudaMemHandler<PixelData<float>, JUST_ALLOC> local_scale_temp2;
442+
ScopedCudaMemHandler<ImgType*, JUST_ALLOC> gradient;
443+
ScopedCudaMemHandler<float*, JUST_ALLOC> local_scale_temp;
444+
ScopedCudaMemHandler<float*, JUST_ALLOC> local_scale_temp2;
444445

445446

446447
// bspline stuff
@@ -490,18 +491,14 @@ class GpuProcessingTask<U>::GpuProcessingTaskImpl {
490491

491492
public:
492493

493-
// TODO: Remove need for passing 'levels' to GpuProcessingTask
494-
// It was used during development to control internal computation like filters, gradient, levels etc. but
495-
// once all is done there is no need for it anymore
496-
GpuProcessingTaskImpl(const PixelData<ImgType> &inputImage, PixelData<float> &levels, const APRParameters &parameters, int maxLevel) :
494+
GpuProcessingTaskImpl(const PixelData<ImgType> &inputImage, const APRParameters &parameters, int maxLevel) :
497495
iCpuImage(inputImage),
498-
iCpuLevels(levels),
499496
iStream(cudaStream.get()),
500497
image (inputImage, iStream),
501498
imageSampling (inputImage, iStream),
502-
gradient (levels, iStream),
503-
local_scale_temp (levels, iStream),
504-
local_scale_temp2 (levels, iStream),
499+
gradient (nullptr, inputImage.getDimensionDS().size(), iStream),
500+
local_scale_temp (nullptr, inputImage.getDimensionDS().size(), iStream),
501+
local_scale_temp2 (nullptr, inputImage.getDimensionDS().size(), iStream),
505502
iParameters(parameters),
506503
iAprInfo(iCpuImage.getDimension()),
507504
iMaxLevel(maxLevel),
@@ -531,7 +528,7 @@ public:
531528
// In LIS we have: var_win[0,1,2] = maximum 3 var_win[3,4,5] = maximum 6
532529
// so maximum paddSize is 6 6 6
533530
PixelDataDim maxPaddSize(6, 6, 6);
534-
PixelDataDim paddedImageSize = levels.getDimension() + maxPaddSize + maxPaddSize;
531+
PixelDataDim paddedImageSize = inputImage.getDimensionDS() + maxPaddSize + maxPaddSize;
535532
lstPadded.initialize(nullptr, paddedImageSize.size(), iStream);
536533
lst2Padded.initialize(nullptr, paddedImageSize.size(), iStream);
537534

@@ -631,22 +628,23 @@ public:
631628
runBsplineOffsetAndCopyOriginal(image.get(), imageSampling.get(), iBsplineOffset /*bspline_offset*/, iCpuImage.getDimension(), iStream);
632629

633630

634-
getGradientCuda(iCpuImage, iCpuLevels, image.get(), gradient.get(), local_scale_temp.get(),
631+
getGradientCuda(iCpuImage, image.get(), gradient.get(), local_scale_temp.get(),
635632
splineCudaX, splineCudaY, splineCudaZ, boundary.get(), isErrorDetectedPinned[0], isErrorDetectedCuda,
636633
iBsplineOffset, iParameters, iStream);
637634

638-
runLocalIntensityScalePipeline(iCpuLevels, iParameters, local_scale_temp.get(), local_scale_temp2.get(), lstPadded.get(), lst2Padded.get(), iStream);
635+
runLocalIntensityScalePipeline(iCpuImage.getDimensionDS(), iParameters, local_scale_temp.get(), local_scale_temp2.get(), lstPadded.get(), lst2Padded.get(), iStream);
639636

640637
// Apply parameters from APRConverter:
641-
runThreshold(local_scale_temp2.get(), gradient.get(), iCpuLevels.x_num, iCpuLevels.y_num, iCpuLevels.z_num, iParameters.Ip_th + iBsplineOffset, iStream);
642-
runRescaleAndThreshold(local_scale_temp.get(), iCpuLevels.mesh.size(), iParameters.sigma_th, iParameters.sigma_th_max, iStream);
643-
runThresholdOpen(gradient.get(), gradient.get(), iCpuLevels.x_num, iCpuLevels.y_num, iCpuLevels.z_num, iParameters.grad_th, iStream);
638+
auto dimOfLevels = iCpuImage.getDimensionDS(); // size of downsampled input image
639+
runThreshold(local_scale_temp2.get(), gradient.get(), dimOfLevels.x, dimOfLevels.y, dimOfLevels.z, iParameters.Ip_th + iBsplineOffset, iStream);
640+
runRescaleAndThreshold(local_scale_temp.get(), dimOfLevels.size(), iParameters.sigma_th, iParameters.sigma_th_max, iStream);
641+
runThresholdOpen(gradient.get(), gradient.get(), dimOfLevels.x, dimOfLevels.y, dimOfLevels.z, iParameters.grad_th, iStream);
644642
// TODO: automatic parameters are not implemented for GPU pipeline (yet)
645643

646644
float min_dim = std::min(iParameters.dy, std::min(iParameters.dx, iParameters.dz));
647645
float level_factor = pow(2, iMaxLevel) * min_dim;
648646
const float mult_const = level_factor/iParameters.rel_error;
649-
runComputeLevels(gradient.get(), local_scale_temp.get(), iCpuLevels.mesh.size(), mult_const, iStream);
647+
runComputeLevels(gradient.get(), local_scale_temp.get(), dimOfLevels.size(), mult_const, iStream);
650648
computeOvpcCuda(local_scale_temp.get(), pctc, iAprInfo, iStream);
651649

652650
computeLinearStructureCuda(y_vec_cuda.get(), xz_end_vec_cuda.get(), level_xz_vec_cuda.get(), pctc, iAprInfo, giga, iParameters, counter_total, iStream);
@@ -660,11 +658,6 @@ public:
660658
// Trim buffer to calculated size (initially it is allocated to worst case - same number of particles as pixels in input image) and copy data from GPU
661659
y_vec.resize(iAprInfo.total_number_particles);
662660
// Copy y_vec from GPU to CPU and synchronize last time - it is needed before we copy data to CPU structures
663-
std::cout << y_vec.size() << "\n";
664-
std::cout << iAprInfo.total_number_particles << "\n";
665-
std::cout << iStream << "\n";
666-
std::cout << y_vec_cuda.getSize() << std::endl;
667-
std::cout << "----------" << std::endl;
668661
checkCuda(cudaMemcpyAsync(y_vec.begin(), y_vec_cuda.get(), iAprInfo.total_number_particles * sizeof(uint16_t), cudaMemcpyDeviceToHost, iStream));
669662

670663

@@ -691,8 +684,8 @@ public:
691684
};
692685

693686
template <typename ImgType>
694-
GpuProcessingTask<ImgType>::GpuProcessingTask(const PixelData<ImgType> &image, PixelData<float> &levels, const APRParameters &parameters, int maxLevel)
695-
: impl{new GpuProcessingTaskImpl<ImgType>(image, levels, parameters, maxLevel)} { }
687+
GpuProcessingTask<ImgType>::GpuProcessingTask(const PixelData<ImgType> &image, const APRParameters &parameters, int maxLevel)
688+
: impl{new GpuProcessingTaskImpl<ImgType>(image, parameters, maxLevel)} { }
696689

697690
template <typename ImgType>
698691
GpuProcessingTask<ImgType>::~GpuProcessingTask() { }
@@ -834,7 +827,7 @@ void getGradient(PixelData<ImgType> &image, PixelData<ImgType> &grad_temp, Pixel
834827
bool isErrorDetected = false;
835828
{
836829
ScopedCudaMemHandler<bool*, JUST_ALLOC> isErrorDetectedCuda(&isErrorDetected, 1, aStream);
837-
getGradientCuda(image, local_scale_temp, cudaImage.get(), cudaGrad.get(), cudalocal_scale_temp.get(),
830+
getGradientCuda(image, cudaImage.get(), cudaGrad.get(), cudalocal_scale_temp.get(),
838831
splineCudaX, splineCudaY, splineCudaZ, boundary.get(), isErrorDetected, isErrorDetectedCuda, bspline_offset, par, aStream);
839832
}
840833
}

‎src/algorithm/ComputeGradientCuda.hpp‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -43,7 +43,7 @@ class GpuProcessingTask {
4343

4444
public:
4545

46-
GpuProcessingTask(const PixelData<ImgType> &image, PixelData<float> &levels, const APRParameters &parameters, int maxLevel);
46+
GpuProcessingTask(const PixelData<ImgType> &image, const APRParameters &parameters, int maxLevel);
4747
~GpuProcessingTask();
4848
GpuProcessingTask(GpuProcessingTask&&);
4949

‎src/algorithm/LocalIntensityScale.cu‎

Lines changed: 7 additions & 7 deletions
Original file line numberDiff line numberDiff line change
@@ -485,12 +485,12 @@ void runConstantScale(S *image, PixelDataDim &dim, cudaStream_t aStream) {
485485
constantScale<<<1, 1, 0, aStream>>>(image, dim.size());
486486
}
487487

488-
template <typename T, typename S>
489-
void runLocalIntensityScalePipeline(const PixelData<T> &image, const APRParameters &par, S *cudaImage, S *cudaTemp, S *lstPadded, S *lst2Padded, cudaStream_t aStream) {
488+
template <typename S>
489+
void runLocalIntensityScalePipeline(const PixelDataDim &tempImageDim, const APRParameters &par, S *cudaImage, S *cudaTemp, S *lstPadded, S *lst2Padded, cudaStream_t aStream) {
490490
float var_rescale;
491491
std::vector<int> var_win;
492492
auto lis = LocalIntensityScale();
493-
lis.get_window_alt(var_rescale, var_win, par, image);
493+
lis.get_window_alt(var_rescale, var_win, par, tempImageDim);
494494
size_t win_y = var_win[0];
495495
size_t win_x = var_win[1];
496496
size_t win_z = var_win[2];
@@ -508,14 +508,14 @@ void runLocalIntensityScalePipeline(const PixelData<T> &image, const APRParamete
508508
constant_scale = true;
509509
}
510510

511-
PixelDataDim imageSize = image.getDimension();
511+
PixelDataDim imageSize = tempImageDim;
512512

513513
if (!constant_scale) {
514514
PixelDataDim paddSize(std::max(win_y, win_y2), std::max(win_x, win_x2), std::max(win_z, win_z2));
515515
PixelDataDim paddedImageSize = imageSize + paddSize + paddSize; // padding on both ends of each dimension
516516
S *ci = cudaImage;
517517
S *ct = cudaTemp;
518-
PixelDataDim dim = image.getDimension();
518+
PixelDataDim dim = tempImageDim;
519519

520520
if (par.reflect_bc_lis) {
521521
runPaddPixels(cudaImage, lstPadded, imageSize, paddedImageSize, paddSize, aStream);
@@ -544,7 +544,7 @@ void runLocalIntensityScalePipeline(const PixelData<T> &image, const APRParamete
544544
}
545545
}
546546

547-
template void runLocalIntensityScalePipeline<float,float>(const PixelData<float>&, const APRParameters&, float*, float*, float*, float*, cudaStream_t);
547+
template void runLocalIntensityScalePipeline<float>(const PixelDataDim &, const APRParameters&, float*, float*, float*, float*, cudaStream_t);
548548

549549

550550

@@ -580,6 +580,6 @@ void getLocalIntensityScale(PixelData<T> &image, PixelData<T> &temp, const APRPa
580580
lstPadded.initialize(nullptr, paddedImageSize.size(), aStream);
581581
lst2Padded.initialize(nullptr, paddedImageSize.size(), aStream);
582582

583-
runLocalIntensityScalePipeline(image, par, cudaImage.get(), cudaTemp.get(), lstPadded.get(), lst2Padded.get(), aStream);
583+
runLocalIntensityScalePipeline(image.getDimension(), par, cudaImage.get(), cudaTemp.get(), lstPadded.get(), lst2Padded.get(), aStream);
584584
}
585585
template void getLocalIntensityScale(PixelData<float>&, PixelData<float>&, const APRParameters&);

‎src/algorithm/LocalIntensityScale.cuh‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -4,7 +4,7 @@
44
#include "data_structures/Mesh/PixelData.hpp"
55
#include "algorithm/APRParameters.hpp"
66

7-
template <typename T, typename S>
8-
void runLocalIntensityScalePipeline(const PixelData<T> &image, const APRParameters &par, S *cudaImage, S *cudaTemp, S *lstPadded, S *lst2Padded, cudaStream_t aStream);
7+
template <typename S>
8+
void runLocalIntensityScalePipeline(const PixelDataDim &image, const APRParameters &par, S *cudaImage, S *cudaTemp, S *lstPadded, S *lst2Padded, cudaStream_t aStream);
99

1010
#endif

‎src/algorithm/LocalIntensityScale.hpp‎

Lines changed: 6 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -36,7 +36,7 @@ void get_local_intensity_scale(PixelData<float> &local_scale_temp, PixelData<flo
3636

3737
float var_rescale;
3838
std::vector<int> var_win;
39-
get_window_alt(var_rescale, var_win, par, local_scale_temp);
39+
get_window_alt(var_rescale, var_win, par, local_scale_temp.getDimension());
4040

4141
int win_y = var_win[0];
4242
int win_x = var_win[1];
@@ -165,8 +165,7 @@ void get_local_intensity_scale(PixelData<float> &local_scale_temp, PixelData<flo
165165

166166
void get_window(float &var_rescale, std::vector<int> &var_win, const APRParameters &par);
167167

168-
template<typename T>
169-
void get_window_alt(float& var_rescale, std::vector<int>& var_win, const APRParameters& par, const PixelData<T>& img);
168+
void get_window_alt(float& var_rescale, std::vector<int>& var_win, const APRParameters& par, const PixelDataDim &img);
170169

171170
template<typename T>
172171
void rescale_var(PixelData<T>& var,const float var_rescale);
@@ -249,8 +248,7 @@ inline void LocalIntensityScale::get_window(float& var_rescale, std::vector<int>
249248
* @param par
250249
* @param temp_img (image already allocated to correct size to compute the local intensity scale)
251250
*/
252-
template<typename T>
253-
inline void LocalIntensityScale::get_window_alt(float& var_rescale, std::vector<int>& var_win, const APRParameters& par,const PixelData<T>& temp_img){
251+
inline void LocalIntensityScale::get_window_alt(float& var_rescale, std::vector<int>& var_win, const APRParameters& par,const PixelDataDim &temp_img){
254252

255253
const double rescale_store_3D[6] = {12.8214, 26.1256, 40.2795, 23.3692, 36.2061, 27.0385};
256254
const double rescale_store_2D[6] = {13.2421, 28.7069, 52.0385, 24.4272, 34.9565, 21.1891};
@@ -267,7 +265,7 @@ inline void LocalIntensityScale::get_window_alt(float& var_rescale, std::vector<
267265

268266
var_win.resize(6,0);
269267

270-
if ( (int) temp_img.y_num > win_val) {
268+
if ( (int) temp_img.y > win_val) {
271269
active_y = true;
272270
var_win[0] = win_1[psf_ind];
273271

@@ -276,15 +274,15 @@ inline void LocalIntensityScale::get_window_alt(float& var_rescale, std::vector<
276274
active_y = false;
277275
}
278276

279-
if ((int) temp_img.x_num > win_val) {
277+
if ((int) temp_img.x > win_val) {
280278
active_x = true;
281279
var_win[1] = win_1[psf_ind];
282280
var_win[4] = win_2[psf_ind];
283281
} else {
284282
active_x = false;
285283
}
286284

287-
if ((int) temp_img.z_num > win_val) {
285+
if ((int) temp_img.z > win_val) {
288286
active_z = true;
289287
var_win[2] = win_1[psf_ind];
290288
var_win[5] = win_2[psf_ind];

‎test/FullPipelineCudaTest.cpp‎

Lines changed: 1 addition & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -303,7 +303,6 @@ namespace {
303303

304304
// Initialize GPU data structures to same values as CPU
305305
PixelData<ImageType> mGpuImage(input_image, true);
306-
PixelData<float> local_scale_temp_GPU(local_scale_temp, false);
307306

308307
// Prepare parameters
309308
APRParameters par;
@@ -340,7 +339,7 @@ namespace {
340339

341340
// Calculate pipeline on GPU
342341
timer.start_timer(">>>>>>>>>>>>>>>>> GPU PIPELINE");
343-
GpuProcessingTask<ImageType> gpt(mGpuImage, local_scale_temp_GPU, par, maxLevel);
342+
GpuProcessingTask<ImageType> gpt(mGpuImage, par, maxLevel);
344343
gpt.processOnGpu();
345344
auto linearAccessGpu = gpt.getDataFromGpu();
346345
giGpu.total_number_particles = linearAccessGpu.y_vec.size();

0 commit comments

Comments
 (0)