diff --git a/CMakeLists.txt b/CMakeLists.txt index 4a10fa2c..985ab1de 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -57,6 +57,8 @@ if(USEGPU) endif() elseif (CUDAToolkit_VERSION_MAJOR STREQUAL "12") set(CUDA_ARCH_LIST "50;52;53;60;61;62;70;72;75;80;86;89;90") + elseif (CUDAToolkit_VERSION_MAJOR STREQUAL "13") + set(CUDA_ARCH_LIST "75;80;86;87;88;89;90;100;103;120;121") else() # some old CUDA version (<10) set(CUDA_ARCH_LIST "50") # 5.0 is the oldest non-deprecated CC as of 2024-09-06 endif() @@ -132,6 +134,9 @@ option(RUN_GTEST "Downloads google unit test API and runs google test scripts to #==== Source files set(SOURCE + # headers + src/nyx/raw_tiff.h + # sources src/nyx/3rdparty/dsyevj3.cpp src/nyx/features/basic_morphology.cpp src/nyx/features/caliper_feret.cpp diff --git a/src/nyx/cache.cpp b/src/nyx/cache.cpp index e1701fb3..4bab7b0f 100644 --- a/src/nyx/cache.cpp +++ b/src/nyx/cache.cpp @@ -9,6 +9,11 @@ return false; \ } \ +#define OKDETL(x,detl) if (x == false) \ +{ \ + return std::string("error due to ") + detl + " at " + __FILE__ + ":" + std::to_string(__LINE__); \ +} \ + #define OKV(x) if (x == false) \ { \ std::cerr << "gpu cache related error in " << __FILE__ << ":" << __LINE__; \ @@ -76,7 +81,7 @@ return szb; } - bool GpusideCache::allocate_gpu_cache( + std::optional GpusideCache::allocate_gpu_cache( // out GpuCache& clouds, // geo moments GpuCache& konturs, @@ -107,14 +112,13 @@ int gabor_ker_side) { using_contour = - using_erosion = - using_gabor = - using_moments = false; + using_erosion = + using_gabor = + using_moments = false; //****** plan GPU memory - size_t amt = 0; - OK(gpu_get_free_mem(amt)); + OKDETL(gpu_get_free_mem(amt), "gpu_get_free_mem()"); int n_gabFilters = n_gabor_filters + 1; // '+1': an extra filter for the baseline signal @@ -130,7 +134,7 @@ n_rois, roi_w, roi_h, n_gabFilters, gabor_ker_side); batch_len = n_rois; - size_t critAmt = amt * 0.75; // 75% GPU RAM as critical RAM + size_t critAmt = float(amt) * 0.75; // 75% GPU RAM as critical RAM if (critAmt < szb) { @@ -139,10 +143,7 @@ { // failed to find a batch ? if (try_nrois == 0) - { - std::cerr << "error: cannot make a ROI batch \n"; - return false; - } + return "cannot rightsize a ROI batch: even 1 ROI requires over critical memory amount " + std::to_string(critAmt); size_t ccl = roi_area * try_nrois; // combined cloud length size_t try_szb = ram_comsumption_szb( @@ -162,7 +163,7 @@ // have we found a compromise ? if (batch_len < n_rois) { - return false; + return "cannot rightsize a ROI batch: batch_len " + std::to_string(batch_len) + " < n_rois " + std::to_string(n_rois); } } @@ -172,8 +173,8 @@ // ROI clouds (always on) - OK(clouds.clear()); - OK(clouds.alloc(batch_roi_cloud_len, batch_len)); + OKDETL(clouds.clear(), "clouds.clear()"); + OKDETL(clouds.alloc(batch_roi_cloud_len, batch_len), "clouds.alloc()"); // contours @@ -181,8 +182,8 @@ { using_contour = true; - OK(konturs.clear()); - OK(konturs.alloc(batch_roi_cloud_len, batch_len)); + OKDETL(konturs.clear(), "konturs.clear()"); + OKDETL(konturs.alloc(batch_roi_cloud_len, batch_len), "konturs.alloc()"); } // moments @@ -192,17 +193,17 @@ using_moments = true; // moments / real intensities - OK(allocate_on_device((void**)realintens, sizeof(RealPixIntens) * batch_roi_cloud_len)); + OKDETL(allocate_on_device((void**)realintens, sizeof(RealPixIntens) * batch_roi_cloud_len), "allocate_on_device(moments1)"); // moments / pre-reduce - OK(allocate_on_device((void**)prereduce, sizeof(double) * batch_roi_cloud_len * 16)); // 16 is the max number of simultaneous totals calculated by a kernel, e.g. RM00-33 + OKDETL(allocate_on_device((void**)prereduce, sizeof(double) * batch_roi_cloud_len * 16), "allocate_on_device(moments2)"); // 16 is the max number of simultaneous totals calculated by a kernel, e.g. RM00-33 // moments / intermediate - OK(intermediate.alloc(GpusideState::__COUNT__, batch_len)); + OKDETL(intermediate.alloc(GpusideState::__COUNT__, batch_len), "allocate_on_device(moments3)"); // moments / CUB DeviceReduce's temp buffer - OK(devicereduce_evaluate_buffer_szb(devicereduce_buf_szb, batch_roi_cloud_len)); - OK(allocate_on_device((void**)devicereduce_buf, devicereduce_buf_szb)); + OKDETL(devicereduce_evaluate_buffer_szb(devicereduce_buf_szb, batch_roi_cloud_len), "allocate_on_device(moments4)"); + OKDETL(allocate_on_device((void**)devicereduce_buf, devicereduce_buf_szb), "allocate_on_device(moments5)"); } // erosion / image matrices 1 and 2 @@ -213,9 +214,9 @@ // imat1 is shared by erosion and Gabor if (!*imat1) - OK(allocate_on_device((void**)imat1, sizeof(imat1[0]) * roi_w * roi_h)); + OKDETL(allocate_on_device((void**)imat1, sizeof(imat1[0]) * roi_w * roi_h), "allocate_on_device(eros1)"); - OK(allocate_on_device((void**)imat2, sizeof(imat2[0]) * roi_w * roi_h)); + OKDETL(allocate_on_device((void**)imat2, sizeof(imat2[0]) * roi_w * roi_h), "allocate_on_device(eros2)"); } // Gabor @@ -226,16 +227,16 @@ // imat1 is shared by erosion and Gabor if (!*imat1) - OK(allocate_on_device((void**)imat1, sizeof(imat1[0]) * roi_w * roi_h)); + OKDETL(allocate_on_device((void**)imat1, sizeof(imat1[0]) * roi_w * roi_h), "allocate_on_device(gabor)"); size_t gabTotlen = (roi_w + gabor_ker_side - 1) * (roi_h + gabor_ker_side - 1) * n_gabFilters; - OK(gabor_linear_image.alloc(gabTotlen, 1)); - OK(gabor_result.alloc(gabTotlen, 1)); - OK(gabor_linear_kernel.alloc(gabTotlen, 1)); - OK(gabor_energy_image.alloc(gabTotlen, 1)); + OKDETL(gabor_linear_image.alloc(gabTotlen, 1), "gabor_linear_image.alloc()"); + OKDETL(gabor_result.alloc(gabTotlen, 1), "gabor_result.alloc()"); + OKDETL(gabor_linear_kernel.alloc(gabTotlen, 1), "gabor_linear_kernel.alloc()"); + OKDETL(gabor_energy_image.alloc(gabTotlen, 1), "gabor_energy_image.alloc()"); } - return true; + return std::nullopt; } void GpusideCache::send_roi_batch_data_2_gpu( diff --git a/src/nyx/cache.h b/src/nyx/cache.h index e129d3dc..0ef11de9 100644 --- a/src/nyx/cache.h +++ b/src/nyx/cache.h @@ -6,6 +6,9 @@ #include "features/pixel.h" #endif +#include +#include + class CpusideCache { public: @@ -148,7 +151,7 @@ class GpusideCache int n_gabor_filters, int gabor_ker_side); - bool allocate_gpu_cache( + std::optional allocate_gpu_cache( // out GpuCache& clouds, GpuCache& konturs, diff --git a/src/nyx/features/contour.cpp b/src/nyx/features/contour.cpp index 39f564ca..86364130 100644 --- a/src/nyx/features/contour.cpp +++ b/src/nyx/features/contour.cpp @@ -583,26 +583,29 @@ ContourFeature::ContourFeature() : FeatureMethod("ContourFeature") // U - unordered subset of C subject to order // S - ordered subset of C - std::list U (C.begin(), C.end()); - + // fix X-crossing points + + std::vector C_fix = C; int n_xxings = 0; - for (Pixel2 p : U) + for (const Pixel2& p : C) { - Pixel2 pN (p.x, p.y-1, p.inten), - pS (p.x, p.y+1, p.inten), - pW (p.x-1, p.y, p.inten), - pE (p.x+1, p.y, p.inten); - auto itN = std::find(U.begin(), U.end(), pN); - auto itS = std::find(U.begin(), U.end(), pS); - auto itE = std::find(U.begin(), U.end(), pE); - auto itW = std::find(U.begin(), U.end(), pW); - if (itN != U.end() && itS != U.end() && itW != U.end() && itE != U.end()) + Pixel2 pN (p.x, p.y - 1, p.inten), + pS (p.x, p.y + 1, p.inten), + pW (p.x - 1, p.y, p.inten), + pE (p.x + 1, p.y, p.inten); + auto itN = std::find (C_fix.begin(), C_fix.end(), pN); + auto itS = std::find (C_fix.begin(), C_fix.end(), pS); + auto itE = std::find (C_fix.begin(), C_fix.end(), pE); + auto itW = std::find (C_fix.begin(), C_fix.end(), pW); + if (itN != C_fix.end() && itS != C_fix.end() && itW != C_fix.end() && itE != C_fix.end()) { - U.remove(p); + auto itP = std::find (C_fix.begin(), C_fix.end(), p); + C_fix.erase (itP); n_xxings++; } } + std::list U (C_fix.begin(), C_fix.end()); // find contour by contour while (!U.empty()) diff --git a/src/nyx/globals.h b/src/nyx/globals.h index 02c4896f..d2fb107b 100644 --- a/src/nyx/globals.h +++ b/src/nyx/globals.h @@ -53,7 +53,6 @@ namespace Nyxus int processDataset_2D_wholeslide( Environment & env, const std::vector& intensFiles, - const std::vector& labelFiles, int n_threads, const SaveOption saveOption, const std::string& outputPath); diff --git a/src/nyx/gpu/erosion.cu b/src/nyx/gpu/erosion.cu index 03fb6858..7169d158 100644 --- a/src/nyx/gpu/erosion.cu +++ b/src/nyx/gpu/erosion.cu @@ -6,7 +6,7 @@ #include #include "gpu.h" -#include "../cache.h" //xxxxxxxxxx #include "../gpucache.h" +#include "../cache.h" #include "../helpers/timing.h" diff --git a/src/nyx/gpu/gabor.cu b/src/nyx/gpu/gabor.cu index 08efed58..612daf90 100644 --- a/src/nyx/gpu/gabor.cu +++ b/src/nyx/gpu/gabor.cu @@ -1,6 +1,6 @@ #include #include "gabor.cuh" -#include "../cache.h" //xxxxxxxxxxxx ../gpucache.h" +#include "../cache.h" using namespace std; diff --git a/src/nyx/gpu/geomoments_central.cu b/src/nyx/gpu/geomoments_central.cu index 84453e6e..753fa1f2 100644 --- a/src/nyx/gpu/geomoments_central.cu +++ b/src/nyx/gpu/geomoments_central.cu @@ -4,7 +4,7 @@ #include #include #include "gpu.h" -#include "../cache.h" //xxxxxxx #include "../gpucache.h" +#include "../cache.h" #include "../features/pixel.h" #include "geomoments.cuh" diff --git a/src/nyx/gpu/geomoments_hu.cu b/src/nyx/gpu/geomoments_hu.cu index 06375c0e..58921f3a 100644 --- a/src/nyx/gpu/geomoments_hu.cu +++ b/src/nyx/gpu/geomoments_hu.cu @@ -7,7 +7,7 @@ #include "../features/image_matrix.h" #include "gpu.h" -#include "../cache.h" //xxxxxxxxx #include "../gpucache.h" +#include "../cache.h" #include "geomoments.cuh" namespace NyxusGpu diff --git a/src/nyx/gpu/geomoments_main.cu b/src/nyx/gpu/geomoments_main.cu index 64dfbcd3..8583195b 100644 --- a/src/nyx/gpu/geomoments_main.cu +++ b/src/nyx/gpu/geomoments_main.cu @@ -7,7 +7,7 @@ #include "gpu.h" #include "geomoments.cuh" -#include "../cache.h" //xxxxxxxxxxxxx #include "../gpucache.h" +#include "../cache.h" #include "../helpers/timing.h" diff --git a/src/nyx/gpu/geomoments_norm.cu b/src/nyx/gpu/geomoments_norm.cu index e19202c8..3a9c1e29 100644 --- a/src/nyx/gpu/geomoments_norm.cu +++ b/src/nyx/gpu/geomoments_norm.cu @@ -7,7 +7,7 @@ #include "../features/image_matrix.h" #include "gpu.h" -#include "../cache.h" //xxxxxxxxxx #include "../gpucache.h" +#include "../cache.h" #include "geomoments.cuh" namespace NyxusGpu diff --git a/src/nyx/gpu/geomoments_origin.cu b/src/nyx/gpu/geomoments_origin.cu index 11ca9e23..a88475b7 100644 --- a/src/nyx/gpu/geomoments_origin.cu +++ b/src/nyx/gpu/geomoments_origin.cu @@ -4,7 +4,7 @@ #include #include #include "gpu.h" -#include "../cache.h" //xxxxx #include "../gpucache.h" +#include "../cache.h" #include "../features/pixel.h" #include "geomoments.cuh" diff --git a/src/nyx/gpu/geomoments_raw.cu b/src/nyx/gpu/geomoments_raw.cu index 57aec400..d59530a9 100644 --- a/src/nyx/gpu/geomoments_raw.cu +++ b/src/nyx/gpu/geomoments_raw.cu @@ -4,7 +4,7 @@ #include #include #include "gpu.h" -#include "../cache.h" //xxxxxxxxx #include "../gpucache.h" +#include "../cache.h" #include "geomoments.cuh" #include "../features/pixel.h" diff --git a/src/nyx/gpu/geomoments_weighting.cu b/src/nyx/gpu/geomoments_weighting.cu index 8590cc8a..3025343e 100644 --- a/src/nyx/gpu/geomoments_weighting.cu +++ b/src/nyx/gpu/geomoments_weighting.cu @@ -113,7 +113,7 @@ namespace NyxusGpu if (pxIdx >= cloud_len) return; - StatsInt x = d_roicloud[pxIdx].x, + [[maybe_unused]] StatsInt x = d_roicloud[pxIdx].x, y = d_roicloud[pxIdx].y; const Pixel2 & p = d_roicloud[pxIdx], @@ -171,4 +171,4 @@ namespace NyxusGpu return true; } -} \ No newline at end of file +} diff --git a/src/nyx/gpu/helper_cuda.h b/src/nyx/gpu/helper_cuda.h index 03ba7400..645fc247 100644 --- a/src/nyx/gpu/helper_cuda.h +++ b/src/nyx/gpu/helper_cuda.h @@ -140,14 +140,14 @@ static const char* _cudaGetErrorEnum(cufftResult error) { case CUFFT_UNALIGNED_DATA: return "CUFFT_UNALIGNED_DATA"; - case CUFFT_INCOMPLETE_PARAMETER_LIST: - return "CUFFT_INCOMPLETE_PARAMETER_LIST"; + //case CUFFT_INCOMPLETE_PARAMETER_LIST: + // return "CUFFT_INCOMPLETE_PARAMETER_LIST"; case CUFFT_INVALID_DEVICE: return "CUFFT_INVALID_DEVICE"; - case CUFFT_PARSE_ERROR: - return "CUFFT_PARSE_ERROR"; + //case CUFFT_PARSE_ERROR: + // return "CUFFT_PARSE_ERROR"; case CUFFT_NO_WORKSPACE: return "CUFFT_NO_WORKSPACE"; @@ -155,8 +155,8 @@ static const char* _cudaGetErrorEnum(cufftResult error) { case CUFFT_NOT_IMPLEMENTED: return "CUFFT_NOT_IMPLEMENTED"; - case CUFFT_LICENSE_ERROR: - return "CUFFT_LICENSE_ERROR"; + //case CUFFT_LICENSE_ERROR: + // return "CUFFT_LICENSE_ERROR"; case CUFFT_NOT_SUPPORTED: return "CUFFT_NOT_SUPPORTED"; @@ -973,4 +973,4 @@ inline bool checkCudaCapabilities(int major_version, int minor_version) { // end of CUDA Helper Functions -#endif // COMMON_HELPER_CUDA_H_ \ No newline at end of file +#endif // COMMON_HELPER_CUDA_H_ diff --git a/src/nyx/gpu/reducers.cu b/src/nyx/gpu/reducers.cu index 91b3defe..5b3b2c71 100644 --- a/src/nyx/gpu/reducers.cu +++ b/src/nyx/gpu/reducers.cu @@ -4,7 +4,7 @@ #include #include #include "gpu.h" -#include "../cache.h" //xxxxxxxxxx #include "../gpucache.h" +#include "../cache.h" #include "../features/pixel.h" #include "geomoments.cuh" diff --git a/src/nyx/gpucache.cpp b/src/nyx/gpucache.cpp index 1c448b77..eabfa3be 100644 --- a/src/nyx/gpucache.cpp +++ b/src/nyx/gpucache.cpp @@ -7,18 +7,6 @@ #include "helpers/helpers.h" #include "roi_cache.h" -#if 0 -// functions implemented in gpucache.cu : -namespace NyxusGpu -{ - bool gpu_delete(void* devptr); - bool allocate_on_device(void** ptr, size_t szb); - bool upload_on_device (void* devbuffer, void* hobuffer, size_t szb); - bool download_on_host(void* hobuffer, void* devbuffer, size_t szb); - bool devicereduce_evaluate_buffer_szb(size_t& devicereduce_buf_szb, size_t maxLen); -} -#endif - #define OK(x) if (x == false) \ { \ std::cerr << "gpu cache related error in " << __FILE__ << ":" << __LINE__; \ @@ -382,444 +370,4 @@ bool GpuCache::download() return true; } -#if 0 -namespace NyxusGpu -{ - size_t ram_comsumption_szb( - bool needContour, - bool needErosion, - bool needGabor, - bool needMoments, - size_t roi_cloud_len, - size_t roi_kontur_cloud_len, - size_t n_rois, - size_t roi_w, - size_t roi_h, - int n_gabor_filters, - int gabor_ker_side) - { - size_t szb_clo = GpuCache::/*clouds.*/szb(roi_cloud_len, n_rois), - szb_state = GpuCache::/*state.*/szb(GpusideState::__COUNT__, n_rois); - - size_t szb_kon = 0; - if (needContour) - { - szb_kon = GpuCache::/*konturs.*/szb(roi_kontur_cloud_len, n_rois); - } - - size_t szb_mom_ri = 0, - szb_moms_pr = 0, - szb_moms_dr = 0; - - if (needMoments) - { - szb_mom_ri = sizeof(RealPixIntens) * roi_cloud_len; - szb_moms_pr = sizeof(double) * roi_cloud_len * 16; - szb_moms_dr = sizeof(void*/*devicereduce_buf[0]*/) * roi_cloud_len; // always enough - } - - size_t szb_imat = 0; - if (needErosion) - { - szb_imat = 2 * sizeof(StatsInt) * roi_w * roi_h; // erosion requires 2 matrices - } - - size_t szb_gabor1 = 0, - szb_gabor2 = 0; - if (needGabor) - { - szb_gabor1 = 3 * sizeof(cufftDoubleComplex) * n_gabor_filters * (roi_w + gabor_ker_side - 1) * (roi_h + gabor_ker_side - 1); // 3 arrays - szb_gabor2 = sizeof(PixIntens) * n_gabor_filters * (roi_w + gabor_ker_side - 1) * (roi_h + gabor_ker_side - 1); - } - - size_t szb = szb_clo + - szb_kon + - szb_mom_ri + - szb_moms_pr + - szb_state + - szb_moms_dr + - szb_imat + - szb_gabor1 + - szb_gabor2; - - return szb; - } - - bool allocate_gpu_cache( - // out - GpuCache& clouds, // geo moments - GpuCache& konturs, - RealPixIntens** realintens, - double** prereduce, - GpuCache& intermediate, - size_t& devicereduce_buf_szb, - void** devicereduce_buf, - size_t & batch_len, - PixIntens** imat1, // erosion - PixIntens** imat2, // (imat1 is shared by erosion and Gabor) - GpuCache & gabor_linear_image, // gabor - GpuCache & gabor_result, - GpuCache & gabor_linear_kernel, - GpuCache & gabor_energy_image, - // in - bool needContour, - bool needErosion, - bool needGabor, - bool needMoments, - size_t roi_cloud_len, - size_t roi_kontur_cloud_len, - size_t n_rois, - size_t roi_area, - size_t roi_w, - size_t roi_h, - int n_gabor_filters, - int gabor_ker_side) - { - NyxusGpu::using_contour = - NyxusGpu::using_erosion = - NyxusGpu::using_gabor = - NyxusGpu::using_moments = false; - - //****** plan GPU memory - - size_t amt = 0; - OK(gpu_get_free_mem(amt)); - - int n_gabFilters = n_gabor_filters + 1; // '+1': an extra filter for the baseline signal - - // Calculate the amt of required memory - size_t ccl0 = roi_area * n_rois; // combined cloud length, initial - size_t szb = ram_comsumption_szb ( - needContour, - needErosion, - needGabor, - needMoments, - ccl0, - roi_kontur_cloud_len, - n_rois, roi_w, roi_h, n_gabFilters, gabor_ker_side); - - batch_len = n_rois; - size_t critAmt = amt * 0.75; // 75% GPU RAM as critical RAM - - if (critAmt < szb) - { - size_t try_nrois = 0; - for (try_nrois = n_rois; try_nrois>=0; try_nrois--) - { - // failed to find a batch ? - if (try_nrois == 0) - { - std::cerr << "error: cannot make a ROI batch \n"; - return false; - } - - size_t ccl = roi_area * try_nrois; // combined cloud length - size_t try_szb = ram_comsumption_szb ( - needContour, - needErosion, - needGabor, - needMoments, - ccl, ccl, try_nrois, roi_w, roi_h, n_gabFilters, gabor_ker_side); - - if (critAmt > try_szb) - { - batch_len = try_nrois; - break; - } - } - - // have we found a compromise ? - if (batch_len < n_rois) - { - std::cerr << "Error: cannot make a ROI batch \n"; - return false; - } - } - - size_t batch_roi_cloud_len = roi_area * batch_len; - - //****** allocate - - // ROI clouds (always on) - - OK(clouds.clear()); - OK(clouds.alloc(batch_roi_cloud_len, batch_len)); - - // contours - - if (needContour) - { - NyxusGpu::using_contour = true; - - OK(konturs.clear()); - OK(konturs.alloc(batch_roi_cloud_len, batch_len)); - } - - // moments - - if (needMoments) - { - using_moments = true; - - // moments / real intensities - OK(NyxusGpu::allocate_on_device((void**)realintens, sizeof(RealPixIntens) * batch_roi_cloud_len)); - - // moments / pre-reduce - OK(NyxusGpu::allocate_on_device((void**)prereduce, sizeof(double) * batch_roi_cloud_len * 16)); // 16 is the max number of simultaneous totals calculated by a kernel, e.g. RM00-33 - - // moments / intermediate - OK(intermediate.alloc(GpusideState::__COUNT__, batch_len)); - - // moments / CUB DeviceReduce's temp buffer - OK(NyxusGpu::devicereduce_evaluate_buffer_szb(devicereduce_buf_szb, batch_roi_cloud_len)); - OK(NyxusGpu::allocate_on_device((void**)devicereduce_buf, devicereduce_buf_szb)); - } - - // erosion / image matrices 1 and 2 - - if (needErosion) - { - using_erosion = true; - - // imat1 is shared by erosion and Gabor - if (! *imat1) - OK(NyxusGpu::allocate_on_device((void**)imat1, sizeof(imat1[0]) * roi_w * roi_h)); - - OK(NyxusGpu::allocate_on_device((void**)imat2, sizeof(imat2[0]) * roi_w * roi_h)); - } - - // Gabor - - if (needGabor) - { - using_gabor = true; - - // imat1 is shared by erosion and Gabor - if (! *imat1) - OK(NyxusGpu::allocate_on_device((void**)imat1, sizeof(imat1[0]) * roi_w * roi_h)); - - size_t gabTotlen = (roi_w + gabor_ker_side - 1) * (roi_h + gabor_ker_side - 1) * n_gabFilters; - OK(gabor_linear_image.alloc(gabTotlen, 1)); - OK(gabor_result.alloc(gabTotlen, 1)); - OK(gabor_linear_kernel.alloc(gabTotlen, 1)); - OK(gabor_energy_image.alloc(gabTotlen, 1)); - } - - return true; - } - - void send_roi_batch_data_2_gpu ( - // out - GpuCache & clouds, - GpuCache & konturs, - RealPixIntens ** realintens, - double ** prereduce, - GpuCache& intermediate, - size_t & devicereduce_buf_szb, - void** devicereduce_buf, - // in - const std::vector& labels, - std::unordered_map & roi_data, - size_t batch_offset, - size_t batch_len) - { - //***** ROI clouds - - // stats of ROI cloud sizes - size_t totCloLen = 0, maxLen = 0; - for (size_t i=batch_offset; i& clouds, - GpuCache& konturs, - RealPixIntens* & realintens, - double* & prereduce, - GpuCache & intermediate, - void* & tempstorage, - PixIntens* imat1, - PixIntens* imat2, - GpuCache & gabor_linear_image, - GpuCache & gabor_result, - GpuCache & gabor_linear_kernel, - GpuCache & gabor_energy_image) - { - // clouds - - OK(clouds.clear()); - - - // contour - - if (using_contour) - { - OK(konturs.clear()); - } - - // moments - - if (using_moments) - { - OK(NyxusGpu::gpu_delete(realintens)); - OK(NyxusGpu::gpu_delete(prereduce)); - OK(intermediate.clear()); - OK(NyxusGpu::gpu_delete(tempstorage)); - realintens = nullptr; - prereduce = nullptr; - OK(intermediate.clear()); - tempstorage = nullptr; // devicereduce's temp storage - } - - // erosion or Gabor - if (using_erosion || using_gabor) - { - OK(NyxusGpu::gpu_delete(imat1)); - } - - // erosion - - if (using_erosion) - { - OK(NyxusGpu::gpu_delete(imat2)); - } - - // Gabor - - if (using_gabor) - { - OK(gabor_linear_image.clear()); - OK(gabor_result.clear()); - OK(gabor_linear_kernel.clear()); - OK(gabor_energy_image.clear()); - } - - return true; - } - -#ifdef USE_GPU - - void send_roi_data_gpuside(const std::vector& roilabels, std::unordered_map & roidata, size_t batch_offset, size_t batch_len) - { - NyxusGpu::send_roi_batch_data_2_gpu( - // out - NyxusGpu::gpu_roiclouds_2d, - NyxusGpu::gpu_roicontours_2d, - &NyxusGpu::dev_realintens, - &NyxusGpu::dev_prereduce, - NyxusGpu::gpu_featurestatebuf, - NyxusGpu::devicereduce_temp_storage_szb, - &NyxusGpu::dev_devicereduce_temp_storage, - // in - roilabels, - roidata, - batch_offset, - batch_len); - } - - // GPU cache of ROI batch data - - GpuCache gpu_roiclouds_2d; - GpuCache gpu_roicontours_2d; - //--later-- GpuCache gpu_roicontours_2d; - RealPixIntens* dev_realintens = nullptr; // max cloud size over batch - double* dev_prereduce = nullptr; // --"-- - GpuCache gpu_featurestatebuf; // n_rois * GeomomentsState::__COUNT__ - - void* dev_devicereduce_temp_storage = nullptr; // allocated [] elements by cub::DeviceReduce::Sum() - size_t devicereduce_temp_storage_szb; - size_t gpu_batch_len = 0; - - // erosion - - PixIntens* dev_imat1; // used by erosion and Gabor - PixIntens* dev_imat2; // erosion only - - // Gabor - - GpuCache gabor_linear_image; - GpuCache gabor_result; - GpuCache gabor_linear_kernel; - GpuCache gabor_energy_image; - -#endif - -} // NyxusGpu -#endif - -#endif +#endif // USE_GPU diff --git a/src/nyx/helpers/helpers.h b/src/nyx/helpers/helpers.h index 9b0f4f7f..55ae1e6c 100644 --- a/src/nyx/helpers/helpers.h +++ b/src/nyx/helpers/helpers.h @@ -48,7 +48,7 @@ namespace Nyxus std::vector L; parse_delimited_string(txt, "\n", L); - size_t maxlen = 0, curlen = 0; + size_t maxlen = 0; for (const auto& l : L) { auto len = l.size(); diff --git a/src/nyx/image_loader.cpp b/src/nyx/image_loader.cpp index 11cf4b0e..f605c88d 100644 --- a/src/nyx/image_loader.cpp +++ b/src/nyx/image_loader.cpp @@ -286,12 +286,12 @@ size_t ImageLoader::get_within_tile_idx (size_t pixel_row, size_t pixel_col) size_t ImageLoader::get_num_tiles_vert() { - return ntw; + return nth; } size_t ImageLoader::get_num_tiles_hor() { - return nth; + return ntw; } size_t ImageLoader::get_tile_height() diff --git a/src/nyx/main_nyxus.cpp b/src/nyx/main_nyxus.cpp index fb314d8d..79be1a0e 100644 --- a/src/nyx/main_nyxus.cpp +++ b/src/nyx/main_nyxus.cpp @@ -85,7 +85,6 @@ int main (int argc, char** argv) errorCode = processDataset_2D_wholeslide( env, intensFiles, - labelFiles, env.n_reduce_threads, env.saveOption, env.output_dir); diff --git a/src/nyx/phase1.cpp b/src/nyx/phase1.cpp index 1bc94940..ea5ead98 100644 --- a/src/nyx/phase1.cpp +++ b/src/nyx/phase1.cpp @@ -13,6 +13,7 @@ #include "environment.h" #include "globals.h" +#include "helpers/helpers.h" #include "helpers/timing.h" namespace Nyxus @@ -29,8 +30,8 @@ namespace Nyxus lyr = 0; // Layer // Read the tiff. The image loader is put in the open state in processDataset() - size_t nth = L.get_num_tiles_hor(), - ntv = L.get_num_tiles_vert(), + size_t ntHor = L.get_num_tiles_hor(), + ntVert = L.get_num_tiles_vert(), fw = L.get_tile_width(), th = L.get_tile_height(), tw = L.get_tile_width(), @@ -38,9 +39,9 @@ namespace Nyxus fullwidth = L.get_full_width(), fullheight = L.get_full_height(); - int cnt = 1; - for (unsigned int row = 0; row < nth; row++) - for (unsigned int col = 0; col < ntv; col++) + size_t tileCnt = 1; + for (size_t row = 0; row < ntVert; row++) + for (size_t col = 0; col < ntHor; col++) { // Fetch the tile bool ok = L.load_tile(row, col); @@ -89,8 +90,8 @@ namespace Nyxus // Show progress info VERBOSLVL2 (env.get_verbosity_level(), - if (cnt++ % 4 == 0) - std::cout << "\t" << int((row * nth + col) * 100 / float(nth * ntv) * 100) / 100. << "%\t" << env.uniqueLabels.size() << " ROIs" << "\n"; + if (tileCnt++ % 4 == 0) + std::cout << "\t" << Nyxus::round2(100. * float(row * ntHor + col) / float(ntHor * ntVert)) << " %\t " << Nyxus::virguler_ulong(env.uniqueLabels.size()) << " ROIs gathered" << "\n"; ); } diff --git a/src/nyx/phase2_25d.cpp b/src/nyx/phase2_25d.cpp index d8711b65..fa06319a 100644 --- a/src/nyx/phase2_25d.cpp +++ b/src/nyx/phase2_25d.cpp @@ -112,7 +112,7 @@ namespace Nyxus float pc = int((row * nth + col) * 100 / float(nth * ntv) * 100) / 100.; if (int(pc) != prevIntPc) { - std::cout << "\t scan trivial " << int(pc) << " %\n"; + std::cout << "\t" << "scan trivial " << int(pc) << " % \n"; prevIntPc = int(pc); } } diff --git a/src/nyx/phase2_2d.cpp b/src/nyx/phase2_2d.cpp index 8f7f7620..8f8e0e45 100644 --- a/src/nyx/phase2_2d.cpp +++ b/src/nyx/phase2_2d.cpp @@ -101,8 +101,8 @@ namespace Nyxus lyr = 0; // Layer // Read the tiffs - size_t nth = ldr.get_num_tiles_hor(), - ntv = ldr.get_num_tiles_vert(), + size_t ntHor = ldr.get_num_tiles_hor(), + ntVert = ldr.get_num_tiles_vert(), fw = ldr.get_tile_width(), th = ldr.get_tile_height(), tw = ldr.get_tile_width(), @@ -110,9 +110,9 @@ namespace Nyxus fullwidth = ldr.get_full_width(), fullheight = ldr.get_full_height(); - int cnt = 1; - for (unsigned int row = 0; row < nth; row++) - for (unsigned int col = 0; col < ntv; col++) + size_t cnt = 1; + for (size_t row = 0; row < ntVert; row++) + for (size_t col = 0; col < ntHor; col++) { // Fetch the tile bool ok = ldr.load_tile(row, col); @@ -162,10 +162,10 @@ namespace Nyxus if (cnt++ % 4 == 0) { static int prevIntPc = 0; - float pc = int((row * nth + col) * 100 / float(nth * ntv) * 100) / 100.; + float pc = Nyxus::round2(100. * float(row * ntHor + col) / float(ntHor * ntVert)); if (int(pc) != prevIntPc) { - std::cout << "\t scan trivial " << int(pc) << " %\n"; + std::cout << "\t" << "scan trivial " << int(pc) << " % \n"; prevIntPc = int(pc); } } @@ -192,8 +192,8 @@ namespace Nyxus lyr = 0; // layer // physical slide properties - size_t nth = ldr.get_num_tiles_hor(), - ntv = ldr.get_num_tiles_vert(), + size_t ntHor = ldr.get_num_tiles_hor(), + ntVert = ldr.get_num_tiles_vert(), fw = ldr.get_tile_width(), th = ldr.get_tile_height(), tw = ldr.get_tile_width(), diff --git a/src/nyx/phase2_3d.cpp b/src/nyx/phase2_3d.cpp index 661b902b..f9444955 100644 --- a/src/nyx/phase2_3d.cpp +++ b/src/nyx/phase2_3d.cpp @@ -40,6 +40,8 @@ namespace Nyxus // time frames: 1:1, 1:N, and N:1 cases are permitted. size_t /* + * we don't need these in the 3D scenario: + * nth = env.theImLoader.get_num_tiles_hor(), ntv = env.theImLoader.get_num_tiles_vert(), fw = env.theImLoader.get_tile_width(), diff --git a/src/nyx/python/new_bindings_py.cpp b/src/nyx/python/new_bindings_py.cpp index 8a377ba3..556dac55 100644 --- a/src/nyx/python/new_bindings_py.cpp +++ b/src/nyx/python/new_bindings_py.cpp @@ -271,7 +271,6 @@ py::tuple featurize_directory_imp( ercode = processDataset_2D_wholeslide( env, intensFiles, - labelFiles, env.n_reduce_threads, env.saveOption, output_path); @@ -631,45 +630,32 @@ py::tuple featurize_fname_lists_imp (uint64_t instid, const py::list& int_fnames // Set the whole-slide/multi-ROI flag theEnvironment.singleROI = single_roi; - std::vector intensFiles, labelFiles; + // Check intensity file names + std::vector intensFiles; for (auto it = int_fnames.begin(); it != int_fnames.end(); ++it) { std::string fn = it->cast(); intensFiles.push_back(fn); } - for (auto it = seg_fnames.begin(); it != seg_fnames.end(); ++it) - { - std::string fn = it->cast(); - labelFiles.push_back(fn); - } - // Check the file names if (intensFiles.size() == 0) - throw std::runtime_error("Intensity file list is blank"); - if (labelFiles.size() == 0) - throw std::runtime_error("Segmentation mask file list is blank"); - if (intensFiles.size() != labelFiles.size()) - throw std::runtime_error("Imbalanced intensity and segmentation mask file lists"); + throw std::runtime_error("Intensity file name list is blank"); + for (auto i = 0; i < intensFiles.size(); i++) { const std::string& i_fname = intensFiles[i]; - const std::string& s_fname = labelFiles[i]; if (!existsOnFilesystem(i_fname)) { auto msg = "File does not exist: " + i_fname; throw std::runtime_error(msg); } - if (!existsOnFilesystem(s_fname)) - { - auto msg = "File does not exist: " + s_fname; - throw std::runtime_error(msg); - } } + // clear result buffers theEnvironment.theResultsCache.clear(); - // Process the image sdata + // Process slides int min_online_roi_size = 0; int errorCode; @@ -681,13 +667,45 @@ py::tuple featurize_fname_lists_imp (uint64_t instid, const py::list& int_fnames } else {return SaveOption::saveBuffer;} }(); - errorCode = processDataset_2D_segmented ( - theEnvironment, - intensFiles, - labelFiles, - theEnvironment.n_reduce_threads, - theEnvironment.saveOption, - output_path); + if (single_roi) + errorCode = processDataset_2D_wholeslide ( + theEnvironment, + intensFiles, + theEnvironment.n_reduce_threads, + theEnvironment.saveOption, + output_path); + else + { + // check mask file names + std::vector labelFiles; + for (auto it = seg_fnames.begin(); it != seg_fnames.end(); ++it) + { + std::string fn = it->cast(); + labelFiles.push_back(fn); + } + + if (intensFiles.size() != labelFiles.size()) + throw std::runtime_error("Imbalanced intensity (" + std::to_string(intensFiles.size()) + " items) and segmentation mask (" + std::to_string(labelFiles.size()) + " items) file lists"); + + for (auto i = 0; i < labelFiles.size(); i++) + { + const std::string& s_fname = labelFiles[i]; + if (!existsOnFilesystem(s_fname)) + { + auto msg = "File does not exist: " + s_fname; + throw std::runtime_error(msg); + } + } + + // we're good to extract features + errorCode = processDataset_2D_segmented( + theEnvironment, + intensFiles, + labelFiles, + theEnvironment.n_reduce_threads, + theEnvironment.saveOption, + output_path); + } if (errorCode) throw std::runtime_error("Error occurred during dataset processing."); diff --git a/src/nyx/python/nyxus/nyxus.py b/src/nyx/python/nyxus/nyxus.py index a2725e07..a189aae3 100644 --- a/src/nyx/python/nyxus/nyxus.py +++ b/src/nyx/python/nyxus/nyxus.py @@ -525,7 +525,7 @@ def featurize_files ( if intensity_files is None: raise IOError ("The list of intensity file paths is empty") - if mask_files is None: + if single_roi == False and mask_files is None: raise IOError ("The list of segment file paths is empty") if (output_type not in self._valid_output_types): diff --git a/src/nyx/raw_image_loader.cpp b/src/nyx/raw_image_loader.cpp index 6c583d56..870e170b 100644 --- a/src/nyx/raw_image_loader.cpp +++ b/src/nyx/raw_image_loader.cpp @@ -152,6 +152,14 @@ bool RawImageLoader::open (const std::string& int_fpath, const std::string& seg_ return true; } +void RawImageLoader::free_tile_buffers() +{ + if (intFL) + intFL->free_tile(); + if (segFL) + segFL->free_tile(); +} + void RawImageLoader::close() { if (segFL) diff --git a/src/nyx/raw_image_loader.h b/src/nyx/raw_image_loader.h index 69ef49af..05d8eb75 100644 --- a/src/nyx/raw_image_loader.h +++ b/src/nyx/raw_image_loader.h @@ -16,6 +16,7 @@ class RawImageLoader void close(); bool load_tile(size_t tile_idx); bool load_tile(size_t tile_row, size_t tile_col); + void free_tile_buffers(); uint32_t get_cur_tile_seg_pixel(size_t pixel_idx); double get_cur_tile_dpequiv_pixel(size_t idx); diff --git a/src/nyx/raw_tiff.h b/src/nyx/raw_tiff.h index e3a2a5bb..b0854e1b 100644 --- a/src/nyx/raw_tiff.h +++ b/src/nyx/raw_tiff.h @@ -172,6 +172,14 @@ class RawTiffTileLoader : public RawFormatLoader // Low level read TIFF bytes auto t_szb = TIFFTileSize(tiff_); tiffTile = _TIFFmalloc(t_szb); + + if (!tiffTile) + { + std::string erm = std::string("_TIFFmalloc() failed at ") + __FILE__ + ":" + std::to_string(__LINE__); + std::cerr << "\n\n" << erm << "\n\n"; + throw std::runtime_error(erm); + } + auto errcode = TIFFReadTile(tiff_, tiffTile, indexColGlobalTile * tileWidth_, indexRowGlobalTile * tileHeight_, 0, 0); if (errcode < 0) { @@ -180,7 +188,7 @@ class RawTiffTileLoader : public RawFormatLoader else // something else { std::string erm = "Tile Loader ERROR: error reading tile data returning code " + std::to_string(errcode); - std::cerr << erm << "\n"; + std::cerr << "\n\n" << erm << "\n\n"; throw std::runtime_error(erm); } } @@ -188,7 +196,8 @@ class RawTiffTileLoader : public RawFormatLoader void free_tile() override { - _TIFFfree(tiffTile); + _TIFFfree (tiffTile); + tiffTile = nullptr; } uint32_t get_uint32_pixel (size_t idx) const diff --git a/src/nyx/slideprops.cpp b/src/nyx/slideprops.cpp index 0881970b..7bb17077 100644 --- a/src/nyx/slideprops.cpp +++ b/src/nyx/slideprops.cpp @@ -197,6 +197,9 @@ namespace Nyxus } } // scan tile + // free tile buffers + ilo.free_tile_buffers(); + #ifdef WITH_PYTHON_H if (PyErr_CheckSignals() != 0) throw pybind11::error_already_set(); diff --git a/src/nyx/workflow_2d_segmented.cpp b/src/nyx/workflow_2d_segmented.cpp index ad747ec0..a6570352 100644 --- a/src/nyx/workflow_2d_segmented.cpp +++ b/src/nyx/workflow_2d_segmented.cpp @@ -61,8 +61,8 @@ namespace Nyxus std::cout << std::setw(15) << freeRamAmt << " b free (" << sgn << memDiff << ") "; ) // Display (1) dataset progress info and (2) file pair info - int digits = std::log10(tot_num_filepairs/100.) + 1, - k = std::pow(10.f, digits); + int digits = std::log10(float(tot_num_filepairs)/100.) + 1, + k = std::pow(10.f, std::abs(digits)); float perCent = float(filepair_index + 1) * 100. / float(tot_num_filepairs); perCent = std::round(perCent * k) / k; VERBOSLVL1 (env.get_verbosity_level(), std::cout << "[ " << filepair_index+1 << " = " << std::setw(digits + 2) << perCent << "% ]\t" << intens_fpath << "\n") @@ -178,7 +178,7 @@ namespace Nyxus size_t nf = intensFiles.size(); { STOPWATCH("prescan/p0/P/#ccbbaa", "\t="); - VERBOSLVL1 (env.get_verbosity_level(), std::cout << "phase 0 (prescanning)\n"); + VERBOSLVL1 (env.get_verbosity_level(), std::cout << "\nphase 0: prescanning " << nf << " slides \n"); env.dataset.reset_dataset_props(); @@ -218,8 +218,7 @@ namespace Nyxus { // allocate VERBOSLVL1 (env.get_verbosity_level(), std::cout << "allocating GPU cache \n"); - - if (! env.devCache.allocate_gpu_cache( + auto allocErr = env.devCache.allocate_gpu_cache( // out env.devCache.gpu_roiclouds_2d, env.devCache.gpu_roicontours_2d, @@ -247,10 +246,10 @@ namespace Nyxus env.dataset.dataset_max_roi_w, env.dataset.dataset_max_roi_h, GaborFeature::f0_theta_pairs.size(), - GaborFeature::n - )) // we need max ROI area inside the function to calculate the batch size if 'dataset_max_combined_roicloud_len' doesn't fit in RAM + GaborFeature::n); + if (allocErr.has_value()) // we need max ROI area inside the function to calculate the batch size if 'dataset_max_combined_roicloud_len' doesn't fit in RAM { - std::cerr << "error in " << __FILE__ << ":" << __LINE__ << "\n"; + std::cerr << "allocating GPU cache failed: " << allocErr.value() << "\n"; return 1; } diff --git a/src/nyx/workflow_2d_whole.cpp b/src/nyx/workflow_2d_whole.cpp index 7db03583..49ebe90f 100644 --- a/src/nyx/workflow_2d_whole.cpp +++ b/src/nyx/workflow_2d_whole.cpp @@ -202,22 +202,26 @@ namespace Nyxus int processDataset_2D_wholeslide ( Environment & env, const std::vector& intensFiles, - const std::vector& labelFiles, int n_threads, const SaveOption saveOption, const std::string& outputPath) { + // create a vector of blank mask file names. Blank mask counterparts + // of intensity files will serve as the condition of the whole-slide scenario in the prescan phase + std::vector labelFiles (intensFiles.size()); + //**** prescan all slides size_t nf = intensFiles.size(); - - VERBOSLVL1 (env.get_verbosity_level(), std::cout << "phase 0 (prescanning)\n"); - + VERBOSLVL1 (env.get_verbosity_level(), std::cout << "\nphase 0: prescanning " << nf << " slides \n"); env.dataset.reset_dataset_props(); for (size_t i=0; i get_3d_compat_phantom() std::string ipath = i_phys_path.string(), mpath = m_phys_path.string(); - //xxxxxxxxxxx return { "C:\\WORK\\AXLE\\OUT\\OUT_synthetic_nifti\\phantom_ngtdm_inten.nii", "C:\\WORK\\AXLE\\OUT\\OUT_synthetic_nifti\\phantom_ngtdm_mask.nii", 57 }; //xxxxxxxxxxxxxxxxxxxxxxxxx - return { ipath, mpath, 1 }; }