From 1ba0a6c266bf9877369922147c48a2abfa6f306c Mon Sep 17 00:00:00 2001 From: Jeremy Kubica <104161096+jeremykubica@users.noreply.github.com> Date: Thu, 1 Feb 2024 11:13:53 -0500 Subject: [PATCH 01/11] Initial commit --- src/kbmod/search/kernels.cu | 62 ++++--- src/kbmod/search/psi_phi_array_utils.h | 33 ---- src/kbmod/search/pydocs/psi_phi_array_docs.h | 95 ++++++---- .../{psi_phi_array.cpp => search_data.cpp} | 163 ++++++++++++------ .../{psi_phi_array_ds.h => search_data_ds.h} | 53 +++--- src/kbmod/search/search_data_utils.h | 39 +++++ src/kbmod/search/stack_search.cpp | 11 +- 7 files changed, 287 insertions(+), 169 deletions(-) delete mode 100644 src/kbmod/search/psi_phi_array_utils.h rename src/kbmod/search/{psi_phi_array.cpp => search_data.cpp} (67%) rename src/kbmod/search/{psi_phi_array_ds.h => search_data_ds.h} (68%) create mode 100644 src/kbmod/search/search_data_utils.h diff --git a/src/kbmod/search/kernels.cu b/src/kbmod/search/kernels.cu index 84b6cf24..2accdb17 100644 --- a/src/kbmod/search/kernels.cu +++ b/src/kbmod/search/kernels.cu @@ -19,29 +19,45 @@ #include "common.h" #include "cuda_errors.h" -#include "psi_phi_array_ds.h" +#include "search_data_ds.h" namespace search { -extern "C" void device_allocate_psi_phi_array(PsiPhiArray *data) { - if (!data->cpu_array_allocated()) throw std::runtime_error("CPU data is not allocated."); - if (data->gpu_array_allocated()) throw std::runtime_error("GPU data is already allocated."); +extern "C" void device_allocate_search_data_arrays(SearchData *data) { + if (!data->cpu_array_allocated() || !data->cpu_time_array_allocated()) { + throw std::runtime_error("CPU data is not allocated."); + } + if (data->gpu_array_allocated() || data->gpu_time_array_allocated()) { + throw std::runtime_error("GPU data is already allocated."); + } + // Allocate space for the psi/phi data. void *device_array_ptr; checkCudaErrors(cudaMalloc((void **)&device_array_ptr, data->get_total_array_size())); checkCudaErrors(cudaMemcpy(device_array_ptr, data->get_cpu_array_ptr(), data->get_total_array_size(), cudaMemcpyHostToDevice)); data->set_gpu_array_ptr(device_array_ptr); + + // Allocate space for the times data. + float *device_times_ptr; + long unsigned time_bytes = data->get_num_times() * sizeof(float); + checkCudaErrors(cudaMalloc((void **)&device_times_ptr, time_bytes)); + checkCudaErrors(cudaMemcpy(device_times_ptr, data->get_cpu_time_array_ptr(), time_bytes, cudaMemcpyHostToDevice)); + data->set_gpu_time_array_ptr(device_times_ptr); } -extern "C" void device_free_psi_phi_array(PsiPhiArray *data) { +extern "C" void device_free_search_data_arrays(SearchData *data) { if (data->gpu_array_allocated()) { checkCudaErrors(cudaFree(data->get_gpu_array_ptr())); data->set_gpu_array_ptr(nullptr); } + if (data->gpu_time_array_allocated()) { + checkCudaErrors(cudaFree(data->get_gpu_time_array_ptr())); + data->set_gpu_time_array_ptr(nullptr); + } } -__forceinline__ __device__ PsiPhi read_encoded_psi_phi(PsiPhiArrayMeta ¶ms, void *psi_phi_vect, int time, +__forceinline__ __device__ PsiPhi read_encoded_psi_phi(SearchDataMeta ¶ms, void *psi_phi_vect, int time, int row, int col) { // Bounds checking. if ((row < 0) || (col < 0) || (row >= params.height) || (col >= params.width)) { @@ -130,7 +146,7 @@ extern "C" __device__ __host__ void SigmaGFilteredIndicesCU(float *values, int n * * Creates a local copy of psi_phi_meta and params in local memory space. */ -__global__ void searchFilterImages(PsiPhiArrayMeta psi_phi_meta, void *psi_phi_vect, float *image_times, +__global__ void searchFilterImages(SearchDataMeta psi_phi_meta, void *psi_phi_vect, float *image_times, SearchParameters params, int num_trajectories, Trajectory *trajectories, Trajectory *results) { // Get the x and y coordinates within the search space. @@ -257,24 +273,25 @@ __global__ void searchFilterImages(PsiPhiArrayMeta psi_phi_meta, void *psi_phi_v } } -extern "C" void deviceSearchFilter(PsiPhiArray &psi_phi_array, float *image_times, SearchParameters params, - int num_trajectories, Trajectory *trj_to_search, int num_results, - Trajectory *best_results) { +extern "C" void deviceSearchFilter(SearchData &search_data, SearchParameters params, int num_trajectories, + Trajectory *trj_to_search, int num_results, Trajectory *best_results) { // Allocate Device memory Trajectory *device_tests; - float *device_img_times; Trajectory *device_search_results; // Check the hard coded maximum number of images against the num_images. - int num_images = psi_phi_array.get_num_times(); + int num_images = search_data.get_num_times(); if (num_images > MAX_NUM_IMAGES) { throw std::runtime_error("Number of images exceeds GPU maximum."); } - // Check that the device psi_phi vector has been allocated. - if (psi_phi_array.gpu_array_allocated() == false) { + // Check that the device vectors have already been allocated. + if (search_data.gpu_array_allocated() == false) { throw std::runtime_error("PsiPhi data has not been created."); } + if (search_data.gpu_time_array_allocated() == false) { + throw std::runtime_error("GPU time data has not been created."); + } // Copy trajectories to search if (params.debug) { @@ -285,14 +302,6 @@ extern "C" void deviceSearchFilter(PsiPhiArray &psi_phi_array, float *image_time checkCudaErrors(cudaMemcpy(device_tests, trj_to_search, sizeof(Trajectory) * num_trajectories, cudaMemcpyHostToDevice)); - // Copy the time vector. - if (params.debug) { - printf("Allocating GPU memory for time data using %lu bytes.\n", sizeof(float) * num_images); - } - checkCudaErrors(cudaMalloc((void **)&device_img_times, sizeof(float) * num_images)); - checkCudaErrors( - cudaMemcpy(device_img_times, image_times, sizeof(float) * num_images, cudaMemcpyHostToDevice)); - // Allocate space for the results. if (params.debug) { printf("Allocating GPU memory for %i results using %lu bytes.\n", num_results, sizeof(Trajectory) * num_results); @@ -308,17 +317,16 @@ extern "C" void deviceSearchFilter(PsiPhiArray &psi_phi_array, float *image_time dim3 threads(THREAD_DIM_X, THREAD_DIM_Y); // Launch Search - searchFilterImages<<>>(psi_phi_array.get_meta_data(), psi_phi_array.get_gpu_array_ptr(), - device_img_times, params, num_trajectories, device_tests, - device_search_results); + searchFilterImages<<>>(search_data.get_meta_data(), search_data.get_gpu_array_ptr(), + static_cast(search_data.get_gpu_time_array_ptr()), params, + num_trajectories, device_tests, device_search_results); // Read back results checkCudaErrors(cudaMemcpy(best_results, device_search_results, sizeof(Trajectory) * num_results, cudaMemcpyDeviceToHost)); - // Free the on GPU memory. + // Free the on GPU memory for this specific search. checkCudaErrors(cudaFree(device_search_results)); - checkCudaErrors(cudaFree(device_img_times)); checkCudaErrors(cudaFree(device_tests)); } diff --git a/src/kbmod/search/psi_phi_array_utils.h b/src/kbmod/search/psi_phi_array_utils.h deleted file mode 100644 index c363f65d..00000000 --- a/src/kbmod/search/psi_phi_array_utils.h +++ /dev/null @@ -1,33 +0,0 @@ -/* - * psi_phi_array_utils.h - * - * The utility functions for the psi/phi array. Broken out from the header - * data structure so that it can use packages that won't be imported into the - * CUDA kernel, such as Eigen. - * - * Created on: Dec 8, 2023 - */ - -#ifndef PSI_PHI_ARRAY_UTILS_ -#define PSI_PHI_ARRAY_UTILS_ - -#include -#include -#include -#include - -#include "common.h" -#include "psi_phi_array_ds.h" -#include "raw_image.h" - -namespace search { - -// Compute the min, max, and scale parameter from the a vector of image data. -std::array compute_scale_params_from_image_vect(const std::vector& imgs, int num_bytes); - -void fill_psi_phi_array(PsiPhiArray& result_data, int num_bytes, const std::vector& psi_imgs, - const std::vector& phi_imgs, bool debug = false); - -} /* namespace search */ - -#endif /* PSI_PHI_ARRAY_UTILS_ */ diff --git a/src/kbmod/search/pydocs/psi_phi_array_docs.h b/src/kbmod/search/pydocs/psi_phi_array_docs.h index ffb7b009..4f668e3e 100644 --- a/src/kbmod/search/pydocs/psi_phi_array_docs.h +++ b/src/kbmod/search/pydocs/psi_phi_array_docs.h @@ -1,5 +1,5 @@ -#ifndef PSI_PHI_ARRAY_DOCS -#define PSI_PHI_ARRAY_DOCS +#ifndef SEARCH_DATA_DOCS +#define SEARCH_DATA_DOCS namespace pydocs { @@ -14,80 +14,88 @@ static const auto DOC_PsiPhi = R"doc( The phi value at a pixel. )doc"; -static const auto DOC_PsiPhiArray = R"doc( +static const auto DOC_SearchData = R"doc( An encoded array of Psi and Phi values along with their meta data. )doc"; -static const auto DOC_PsiPhiArray_get_num_bytes = R"doc( +static const auto DOC_SearchData_get_num_bytes = R"doc( The target number of bytes to use for encoding the data (1 for uint8, 2 for uint16, or 4 for float32). Might differ from actual number of bytes (block_size). )doc"; -static const auto DOC_PsiPhiArray_get_num_times = R"doc( +static const auto DOC_SearchData_get_num_times = R"doc( The number of times. )doc"; -static const auto DOC_PsiPhiArray_get_width = R"doc( +static const auto DOC_SearchData_get_width = R"doc( The image width. )doc"; -static const auto DOC_PsiPhiArray_get_height = R"doc( +static const auto DOC_SearchData_get_height = R"doc( The image height. )doc"; -static const auto DOC_PsiPhiArray_get_pixels_per_image = R"doc( +static const auto DOC_SearchData_get_pixels_per_image = R"doc( The number of pixels per each image. )doc"; -static const auto DOC_PsiPhiArray_get_num_entries = R"doc( +static const auto DOC_SearchData_get_num_entries = R"doc( The number of array entries. )doc"; -static const auto DOC_PsiPhiArray_get_total_array_size = R"doc( +static const auto DOC_SearchData_get_total_array_size = R"doc( The size of the array in bytes. )doc"; -static const auto DOC_PsiPhiArray_get_block_size = R"doc( +static const auto DOC_SearchData_get_block_size = R"doc( The size of a single entry in bytes. )doc"; -static const auto DOC_PsiPhiArray_get_psi_min_val = R"doc( +static const auto DOC_SearchData_get_psi_min_val = R"doc( The minimum value of psi used in the scaling computations. )doc"; -static const auto DOC_PsiPhiArray_get_psi_max_val = R"doc( +static const auto DOC_SearchData_get_psi_max_val = R"doc( The maximum value of psi used in the scaling computations. )doc"; -static const auto DOC_PsiPhiArray_get_psi_scale = R"doc( +static const auto DOC_SearchData_get_psi_scale = R"doc( The scaling parameter for psi. )doc"; -static const auto DOC_PsiPhiArray_get_phi_min_val = R"doc( +static const auto DOC_SearchData_get_phi_min_val = R"doc( The minimum value of phi used in the scaling computations. )doc"; -static const auto DOC_PsiPhiArray_get_phi_max_val = R"doc( +static const auto DOC_SearchData_get_phi_max_val = R"doc( The maximum value of phi used in the scaling computations. )doc"; -static const auto DOC_PsiPhiArray_get_phi_scale = R"doc( +static const auto DOC_SearchData_get_phi_scale = R"doc( The scaling parameter for phi. )doc"; -static const auto DOC_PsiPhiArray_get_cpu_array_allocated = R"doc( - A Boolean indicating whether the cpu array exists. +static const auto DOC_SearchData_get_cpu_array_allocated = R"doc( + A Boolean indicating whether the cpu data (psi/phi) array exists. )doc"; -static const auto DOC_PsiPhiArray_get_gpu_array_allocated = R"doc( - A Boolean indicating whether the gpu array exists. +static const auto DOC_SearchData_get_gpu_array_allocated = R"doc( + A Boolean indicating whether the gpu data (psi/phi) array exists. )doc"; -static const auto DOC_PsiPhiArray_clear = R"doc( +static const auto DOC_SearchData_get_cpu_time_array_allocated = R"doc( + A Boolean indicating whether the cpu time array exists. + )doc"; + +static const auto DOC_SearchData_get_gpu_time_array_allocated = R"doc( + A Boolean indicating whether the gpu time array exists. + )doc"; + +static const auto DOC_SearchData_clear = R"doc( Clear all data and free the arrays. )doc"; -static const auto DOC_PsiPhiArray_read_psi_phi = R"doc( +static const auto DOC_SearchData_read_psi_phi = R"doc( Read a PsiPhi value from the CPU array. Parameters @@ -105,9 +113,23 @@ static const auto DOC_PsiPhiArray_read_psi_phi = R"doc( The pixel values. )doc"; -static const auto DOC_PsiPhiArray_set_meta_data = R"doc( +static const auto DOC_SearchData_read_time = R"doc( + Read a zeroed time value from the CPU array. + + Parameters + ---------- + time : `int` + The timestep to read. + + Returns + ------- + `float` + The time. + )doc"; + +static const auto DOC_SearchData_set_meta_data = R"doc( Set the meta data for the array. Automatically called by - fill_psi_phi_array(). + fill_search_data(). Parameters ---------- @@ -121,12 +143,12 @@ static const auto DOC_PsiPhiArray_set_meta_data = R"doc( The width of each image in pixels. )doc"; -static const auto DOC_PsiPhiArray_fill_psi_phi_array = R"doc( - Fill the PsiPhiArray from Psi and Phi images. +static const auto DOC_SearchData_fill_search_data = R"doc( + Fill the SearchData from Psi and Phi images. Parameters ---------- - result_data : `PsiPhiArray` + result_data : `SearchData` The location to store the data. num_bytes : `int` The type of encoding to use (1, 2, or 4). @@ -134,8 +156,23 @@ static const auto DOC_PsiPhiArray_fill_psi_phi_array = R"doc( A list of psi images. phi_imgs : `list` A list of phi images. + zeroed_times : `list` + A list of floating point times starting at zero. + )doc"; + +static const auto DOC_SearchData_fill_search_data_from_image_stack = R"doc( + Fill the SearchData an ImageStack. + + Parameters + ---------- + result_data : `SearchData` + The location to store the data. + num_bytes : `int` + The type of encoding to use (1, 2, or 4). + stack : `ImageStack` + The stack of LayeredImages from which to build the psi and phi images. )doc"; } // namespace pydocs -#endif /* PSI_PHI_ARRAY_DOCS */ +#endif /* SEARCH_DATA_DOCS */ diff --git a/src/kbmod/search/psi_phi_array.cpp b/src/kbmod/search/search_data.cpp similarity index 67% rename from src/kbmod/search/psi_phi_array.cpp rename to src/kbmod/search/search_data.cpp index 3f265668..4331f779 100644 --- a/src/kbmod/search/psi_phi_array.cpp +++ b/src/kbmod/search/search_data.cpp @@ -1,43 +1,41 @@ -#include "psi_phi_array_ds.h" -#include "psi_phi_array_utils.h" -#include "pydocs/psi_phi_array_docs.h" +#include "search_data_ds.h" +#include "search_data_utils.h" +#include "pydocs/search_data_docs.h" namespace search { // Declaration of CUDA functions that will be linked in. #ifdef HAVE_CUDA -extern "C" void device_allocate_psi_phi_array(PsiPhiArray* data); +extern "C" void device_allocate_search_data_arrays(SearchData* data); -extern "C" void device_free_psi_phi_array(PsiPhiArray* data); +extern "C" void device_free_search_data_arrays(SearchData* data); #endif // ------------------------------------------------------- // --- Implementation of core data structure functions --- // ------------------------------------------------------- -PsiPhiArray::PsiPhiArray() {} +SearchData::SearchData() {} -PsiPhiArray::~PsiPhiArray() { - if (cpu_array_ptr != nullptr) { - free(cpu_array_ptr); - } -#ifdef HAVE_CUDA - if (gpu_array_ptr != nullptr) { - device_free_psi_phi_array(this); - } -#endif +SearchData::~SearchData() { + clear(); } -void PsiPhiArray::clear() { +void SearchData::clear() { // Free all used memory on CPU and GPU. if (cpu_array_ptr != nullptr) { free(cpu_array_ptr); cpu_array_ptr = nullptr; } + if (cpu_time_array != nullptr) { + free(cpu_time_array); + cpu_array_ptr = nullptr; + } #ifdef HAVE_CUDA - if (gpu_array_ptr != nullptr) { - device_free_psi_phi_array(this); + if ((gpu_array_ptr != nullptr) || (gpu_time_array != nullptr)) { + device_free_search_data_arrays(this); gpu_array_ptr = nullptr; + gpu_time_array = nullptr; } #endif @@ -58,7 +56,7 @@ void PsiPhiArray::clear() { meta_data.phi_scale = 1.0; } -void PsiPhiArray::set_meta_data(int new_num_bytes, int new_num_times, int new_height, int new_width) { +void SearchData::set_meta_data(int new_num_bytes, int new_num_times, int new_height, int new_width) { // Validity checking of parameters. if (new_num_bytes != -1 && new_num_bytes != 1 && new_num_bytes != 2 && new_num_bytes != 4) { throw std::runtime_error("Invalid setting of num_bytes. Must be (-1 [use default], 1, 2, or 4)."); @@ -90,7 +88,7 @@ void PsiPhiArray::set_meta_data(int new_num_bytes, int new_num_times, int new_he meta_data.total_array_size = meta_data.block_size * meta_data.num_entries; } -void PsiPhiArray::set_psi_scaling(float min_val, float max_val, float scale_val) { +void SearchData::set_psi_scaling(float min_val, float max_val, float scale_val) { if (min_val > max_val) throw std::runtime_error("Min value needs to be < max value"); if (scale_val <= 0) throw std::runtime_error("Scale value must be greater than zero."); meta_data.psi_min_val = min_val; @@ -98,7 +96,7 @@ void PsiPhiArray::set_psi_scaling(float min_val, float max_val, float scale_val) meta_data.psi_scale = scale_val; } -void PsiPhiArray::set_phi_scaling(float min_val, float max_val, float scale_val) { +void SearchData::set_phi_scaling(float min_val, float max_val, float scale_val) { if (min_val > max_val) throw std::runtime_error("Min value needs to be < max value"); if (scale_val <= 0) throw std::runtime_error("Scale value must be greater than zero."); meta_data.phi_min_val = min_val; @@ -106,7 +104,7 @@ void PsiPhiArray::set_phi_scaling(float min_val, float max_val, float scale_val) meta_data.phi_scale = scale_val; } -PsiPhi PsiPhiArray::read_psi_phi(int time, int row, int col) { +PsiPhi SearchData::read_psi_phi(int time, int row, int col) { PsiPhi result = {NO_DATA, NO_DATA}; // Array allocation and bounds checking. @@ -140,6 +138,15 @@ PsiPhi PsiPhiArray::read_psi_phi(int time, int row, int col) { return result; } +float SearchData::read_time(int time_index) { + if (cpu_time_array == nullptr) throw std::runtime_error("Read from unallocated times array."); + if ((time_index < 0 )|| (time_index >= meta_data.num_times)) { + throw std::runtime_error("Out of bounds read for time step=%i", time_index); + } + return cpu_time_array[time_index]; +} + + // ------------------------------------------- // --- Implementation of utility functions --- // ------------------------------------------- @@ -171,7 +178,7 @@ std::array compute_scale_params_from_image_vect(const std::vector -void set_encode_cpu_psi_phi_array(PsiPhiArray& data, const std::vector& psi_imgs, +void set_encode_cpu_search_data(SearchData& data, const std::vector& psi_imgs, const std::vector& phi_imgs, bool debug) { if (data.get_cpu_array_ptr() != nullptr) { throw std::runtime_error("CPU PsiPhi already allocated."); @@ -215,7 +222,7 @@ void set_encode_cpu_psi_phi_array(PsiPhiArray& data, const std::vector data.set_cpu_array_ptr((void*)encoded); } -void set_float_cpu_psi_phi_array(PsiPhiArray& data, const std::vector& psi_imgs, +void set_float_cpu_search_data(SearchData& data, const std::vector& psi_imgs, const std::vector& phi_imgs, bool debug) { if (data.get_cpu_array_ptr() != nullptr) { throw std::runtime_error("CPU PsiPhi already allocated."); @@ -241,8 +248,9 @@ void set_float_cpu_psi_phi_array(PsiPhiArray& data, const std::vector& data.set_cpu_array_ptr((void*)encoded); } -void fill_psi_phi_array(PsiPhiArray& result_data, int num_bytes, const std::vector& psi_imgs, - const std::vector& phi_imgs, bool debug) { +void fill_search_data(SearchData& result_data, int num_bytes, const std::vector& psi_imgs, + const std::vector& phi_imgs, const std::vector zeroed_times, + bool debug) { if (result_data.get_cpu_array_ptr() != nullptr) { return; } @@ -251,6 +259,7 @@ void fill_psi_phi_array(PsiPhiArray& result_data, int num_bytes, const std::vect int num_times = psi_imgs.size(); if (num_times <= 0) throw std::runtime_error("Trying to fill PsiPhi from empty vectors."); if (num_times != phi_imgs.size()) throw std::runtime_error("Size mismatch between psi and phi."); + if (num_times != zeroed_times.size()) throw std::runtime_error("Size mismatch between psi and zeroed times."); int width = phi_imgs[0].get_width(); int height = phi_imgs[0].get_height(); @@ -275,78 +284,126 @@ void fill_psi_phi_array(PsiPhiArray& result_data, int num_bytes, const std::vect // Do the local encoding. if (result_data.get_num_bytes() == 1) { - set_encode_cpu_psi_phi_array(result_data, psi_imgs, phi_imgs, debug); + set_encode_cpu_search_data(result_data, psi_imgs, phi_imgs, debug); } else { - set_encode_cpu_psi_phi_array(result_data, psi_imgs, phi_imgs, debug); + set_encode_cpu_search_data(result_data, psi_imgs, phi_imgs, debug); } } else { if (debug) { printf("Encoding psi and phi as floats.\n"); } // Just interleave psi and phi images. - set_float_cpu_psi_phi_array(result_data, psi_imgs, phi_imgs, debug); + set_float_cpu_search_data(result_data, psi_imgs, phi_imgs, debug); } + // Copy the time array. + const long unsigned times_bytes = result_data.get_num_times() * sizeof(float); + if (debug) printf("Allocating %lu bytes on the CPU for times.\n"); + + float* times_array = (float*)malloc(data.get_total_array_size()); + if (times_array == nullptr) throw std::runtime_error("Unable to allocate space for CPU times."); + for (int i = 0; i < result_data.get_num_times(); ++i) { + times_array[i] = zeroed_times[i]; + } + data.set_cpu_time_array_ptr((void*)times_array); + #ifdef HAVE_CUDA // Create a copy of the encoded data in GPU memory. if (debug) { printf("Allocating GPU memory for PsiPhi array using %lu bytes.\n", result_data.get_total_array_size()); + printf("Allocating GPU memory for times array using %lu bytes.\n", times_bytes); } - device_allocate_psi_phi_array(&result_data); + + device_allocate_search_data_arrays(&result_data); if (result_data.get_gpu_array_ptr() == nullptr) { throw std::runtime_error("Unable to allocate GPU PsiPhi array."); } + if (result_data.get_gpu_time_array_ptr() == nullptr) { + throw std::runtime_error("Unable to allocate GPU time array."); + } #endif } +void fill_search_data_from_image_stack(SearchData& result_data, ImageStack& stack, int num_bytes, bool debug) { + // Compute Phi and Psi from convolved images while leaving masked pixels alone + // Reinsert 0s for NO_DATA? + std::vector psi_images; + std::vector phi_images; + const int num_images = stack.img_count(); + if (debug) { + unsigned long num_bytes = 2 * stack.get_height() * stack.get_width() * num_images * sizeof(float); + printf("Building %i temporary %i by %i images (psi and phi), requiring %lu bytes", + (num_images * 2), stack.get_width(), stack.get_height(), num_bytes); + } + + // Build the psi and phi images first. + for (int i = 0; i < num_images; ++i) { + LayeredImage& img = stack.get_single_image(i); + psi_images.push_back(img.generate_psi_image()); + phi_images.push_back(img.generate_phi_image()); + } + + // Convert these into an array form. Needs the full psi and phi computed first so the + // encoding can compute the bounds of each array. + std::vector zeroed_times = stack.build_zeroed_times(); + fill_search_data(psi_phi_data, num_bytes, psi_images, phi_images, zeroed_times, debug); +} + // ------------------------------------------- // --- Python definitions -------------------- // ------------------------------------------- #ifdef Py_PYTHON_H -static void psi_phi_array_binding(py::module& m) { - using ppa = search::PsiPhiArray; +static void search_data_binding(py::module& m) { + using ppa = search::SearchData; py::class_(m, "PsiPhi", pydocs::DOC_PsiPhi) .def(py::init<>()) .def_readwrite("psi", &search::PsiPhi::psi) .def_readwrite("phi", &search::PsiPhi::phi); - py::class_(m, "PsiPhiArray", pydocs::DOC_PsiPhiArray) + py::class_(m, "SearchData", pydocs::DOC_SearchData) .def(py::init<>()) - .def_property_readonly("num_bytes", &ppa::get_num_bytes, pydocs::DOC_PsiPhiArray_get_num_bytes) - .def_property_readonly("num_times", &ppa::get_num_times, pydocs::DOC_PsiPhiArray_get_num_times) - .def_property_readonly("width", &ppa::get_width, pydocs::DOC_PsiPhiArray_get_width) - .def_property_readonly("height", &ppa::get_height, pydocs::DOC_PsiPhiArray_get_height) + .def_property_readonly("num_bytes", &ppa::get_num_bytes, pydocs::DOC_SearchData_get_num_bytes) + .def_property_readonly("num_times", &ppa::get_num_times, pydocs::DOC_SearchData_get_num_times) + .def_property_readonly("width", &ppa::get_width, pydocs::DOC_SearchData_get_width) + .def_property_readonly("height", &ppa::get_height, pydocs::DOC_SearchData_get_height) .def_property_readonly("pixels_per_image", &ppa::get_pixels_per_image, - pydocs::DOC_PsiPhiArray_get_pixels_per_image) + pydocs::DOC_SearchData_get_pixels_per_image) .def_property_readonly("num_entries", &ppa::get_num_entries, - pydocs::DOC_PsiPhiArray_get_num_entries) + pydocs::DOC_SearchData_get_num_entries) .def_property_readonly("total_array_size", &ppa::get_total_array_size, - pydocs::DOC_PsiPhiArray_get_total_array_size) - .def_property_readonly("block_size", &ppa::get_block_size, pydocs::DOC_PsiPhiArray_get_block_size) + pydocs::DOC_SearchData_get_total_array_size) + .def_property_readonly("block_size", &ppa::get_block_size, pydocs::DOC_SearchData_get_block_size) .def_property_readonly("psi_min_val", &ppa::get_psi_min_val, - pydocs::DOC_PsiPhiArray_get_psi_min_val) + pydocs::DOC_SearchData_get_psi_min_val) .def_property_readonly("psi_max_val", &ppa::get_psi_max_val, - pydocs::DOC_PsiPhiArray_get_psi_max_val) - .def_property_readonly("psi_scale", &ppa::get_psi_scale, pydocs::DOC_PsiPhiArray_get_psi_scale) + pydocs::DOC_SearchData_get_psi_max_val) + .def_property_readonly("psi_scale", &ppa::get_psi_scale, pydocs::DOC_SearchData_get_psi_scale) .def_property_readonly("phi_min_val", &ppa::get_phi_min_val, - pydocs::DOC_PsiPhiArray_get_phi_min_val) + pydocs::DOC_SearchData_get_phi_min_val) .def_property_readonly("phi_max_val", &ppa::get_phi_max_val, - pydocs::DOC_PsiPhiArray_get_phi_max_val) - .def_property_readonly("phi_scale", &ppa::get_phi_scale, pydocs::DOC_PsiPhiArray_get_phi_scale) + pydocs::DOC_SearchData_get_phi_max_val) + .def_property_readonly("phi_scale", &ppa::get_phi_scale, pydocs::DOC_SearchData_get_phi_scale) .def_property_readonly("cpu_array_allocated", &ppa::cpu_array_allocated, - pydocs::DOC_PsiPhiArray_get_cpu_array_allocated) + pydocs::DOC_SearchData_get_cpu_array_allocated) .def_property_readonly("gpu_array_allocated", &ppa::gpu_array_allocated, - pydocs::DOC_PsiPhiArray_get_gpu_array_allocated) - .def("set_meta_data", &ppa::set_meta_data, pydocs::DOC_PsiPhiArray_set_meta_data) - .def("clear", &ppa::clear, pydocs::DOC_PsiPhiArray_clear) - .def("read_psi_phi", &ppa::read_psi_phi, pydocs::DOC_PsiPhiArray_read_psi_phi); + pydocs::DOC_SearchData_get_gpu_array_allocated) + .def_property_readonly("cpu_time_array_allocated", &ppa::cpu_time_array_allocated, + pydocs::DOC_SearchData_get_cpu_time_array_allocated) + .def_property_readonly("gpu_time_array_allocated", &ppa::gpu_time_array_allocated, + pydocs::DOC_SearchData_get_gpu_time_array_allocated) + .def("set_meta_data", &ppa::set_meta_data, pydocs::DOC_SearchData_set_meta_data) + .def("clear", &ppa::clear, pydocs::DOC_SearchData_clear) + .def("read_psi_phi", &ppa::read_psi_phi, pydocs::DOC_SearchData_read_psi_phi); + .def("read_time", &ppa::read_time, pydocs::DOC_SearchData_read_time); m.def("compute_scale_params_from_image_vect", &search::compute_scale_params_from_image_vect); m.def("decode_uint_scalar", &search::decode_uint_scalar); m.def("encode_uint_scalar", &search::encode_uint_scalar); - m.def("fill_psi_phi_array", &search::fill_psi_phi_array, pydocs::DOC_PsiPhiArray_fill_psi_phi_array); + m.def("fill_search_data", &search::fill_search_data, pydocs::DOC_SearchData_fill_search_data); + m.def("fill_search_data_from_image_stack", &search::fill_search_data_from_image_stack, + pydocs::DOC_SearchData_fill_search_data_from_image_stack); } #endif diff --git a/src/kbmod/search/psi_phi_array_ds.h b/src/kbmod/search/search_data_ds.h similarity index 68% rename from src/kbmod/search/psi_phi_array_ds.h rename to src/kbmod/search/search_data_ds.h index 719ec868..900c3cf5 100644 --- a/src/kbmod/search/psi_phi_array_ds.h +++ b/src/kbmod/search/search_data_ds.h @@ -1,22 +1,23 @@ /* - * psi_phi_array_ds.h + * search_data_ds.h * - * The data structure for the interleaved psi/phi array. The the data + * The data structure for the raw data needed for the search algorith, + * including the psi/phi values and the zeroed times. The the data * structure and core functions are included in the header (and separated out * from the rest of the utility functions) to allow the CUDA files to import * only what they need. * * The data structure allocates memory on both the CPU and GPU for the - * interleaved psi/phi array and maintains ownership of the pointers - * until clear() is called or the PsiPhiArray's destructor is called. This allows - * the object to be passed repeatedly to the on-device search without reallocating - * and copying the memory on the GPU. + * arraysand maintains ownership of the pointers until clear() is called + * the object's destructor is called. This allows the object to be passed + * repeatedly to the on-device search without reallocating and copying the + * memory on the GPU. * * Created on: Dec 5, 2023 */ -#ifndef PSI_PHI_ARRAY_DS_ -#define PSI_PHI_ARRAY_DS_ +#ifndef SEARCH_DATA_DS_ +#define SEARCH_DATA_DS_ #include #include @@ -42,8 +43,8 @@ inline float decode_uint_scalar(float value, float min_val, float scale) { return (value == 0.0) ? NO_DATA : (value - 1.0) * scale + min_val; } -// The struct of meta data for the PsiPhiArray. -struct PsiPhiArrayMeta { +// The struct of meta data for the SearchData. +struct SearchDataMeta { int num_times = 0; int width = 0; int height = 0; @@ -64,17 +65,17 @@ struct PsiPhiArrayMeta { float phi_scale = 1.0; }; -/* PsiPhiArray is a class to hold the psi and phi arrays for the CPU and GPU as well as +/* SearchData is a class to hold the psi and phi arrays for the CPU and GPU as well as the meta data and functions to do encoding and decoding on CPU. */ -class PsiPhiArray { +class SearchData { public: - explicit PsiPhiArray(); - virtual ~PsiPhiArray(); + explicit SearchData(); + virtual ~SearchData(); void clear(); - inline PsiPhiArrayMeta& get_meta_data() { return meta_data; } + inline SearchDataMeta& get_meta_data() { return meta_data; } // --- Getter functions (for Python interface) ---------------- inline int get_num_bytes() { return meta_data.num_bytes; } @@ -95,9 +96,12 @@ class PsiPhiArray { inline bool cpu_array_allocated() { return cpu_array_ptr != nullptr; } inline bool gpu_array_allocated() { return gpu_array_ptr != nullptr; } + inline bool cpu_time_array_allocated() { return cpu_time_array != nullptr; } + inline bool gpu_time_array_allocated() { return gpu_time_array != nullptr; } - // Primary getter function for interaction (read the data). - PsiPhi read_psi_phi(int time, int row, int col); + // Primary getter functions for interaction (read the data). + PsiPhi read_psi_phi(int time_index, int row, int col); + float read_time_value(int time_index); // Setters for the utility functions to allocate the data. void set_meta_data(int new_num_bytes, int new_num_times, int new_height, int new_width); @@ -110,14 +114,21 @@ class PsiPhiArray { inline void set_cpu_array_ptr(void* new_ptr) { cpu_array_ptr = new_ptr; } inline void set_gpu_array_ptr(void* new_ptr) { gpu_array_ptr = new_ptr; } + inline float* get_cpu_time_array_ptr() { return cpu_time_array; } + inline float* get_gpu_time_array_ptr() { return gpu_time_array; } + inline void set_cpu_time_array_ptr(float* new_ptr) { cpu_time_array = new_ptr; } + inline void set_gpu_time_array_ptr(float* new_ptr) { gpu_time_array = new_ptr; } + private: - PsiPhiArrayMeta meta_data; + SearchDataMeta meta_data; - // Pointers the array (CPU space). - void* cpu_array_ptr = nullptr; + // Pointers to the arrays + void* cpu_array_ptr = nullptr; void* gpu_array_ptr = nullptr; + float* cpu_time_array = nullptr; + float* gpu_time_array = nullptr; }; } /* namespace search */ -#endif /* PSI_PHI_ARRAY_DS_ */ +#endif /* SEARCH_DATA_DS_ */ diff --git a/src/kbmod/search/search_data_utils.h b/src/kbmod/search/search_data_utils.h new file mode 100644 index 00000000..50538e99 --- /dev/null +++ b/src/kbmod/search/search_data_utils.h @@ -0,0 +1,39 @@ +/* + * search_data_utils.h + * + * The utility functions for the psi/phi array. Broken out from the header + * data structure so that it can use packages that won't be imported into the + * CUDA kernel, such as Eigen. + * + * Created on: Dec 8, 2023 + */ + +#ifndef SEARCH_DATA_UTILS_ +#define SEARCH_DATA_UTILS_ + +#include +#include +#include +#include + +#include "common.h" +#include "image_stack.h" +#include "layered_image.h" +#include "search_data_ds.h" +#include "raw_image.h" + +namespace search { + +// Compute the min, max, and scale parameter from the a vector of image data. +std::array compute_scale_params_from_image_vect(const std::vector& imgs, int num_bytes); + +void fill_search_data(SearchData& result_data, int num_bytes, const std::vector& psi_imgs, + const std::vector& phi_imgs, const std::vector zeroed_times, + bool debug = false); + +void fill_search_data_from_image_stack(SearchData& result_data, ImageStack& stack, int num_bytes, + bool debug = false); + +} /* namespace search */ + +#endif /* SEARCH_DATA_UTILS_ */ diff --git a/src/kbmod/search/stack_search.cpp b/src/kbmod/search/stack_search.cpp index ba31ef24..60ff1611 100644 --- a/src/kbmod/search/stack_search.cpp +++ b/src/kbmod/search/stack_search.cpp @@ -2,9 +2,8 @@ namespace search { #ifdef HAVE_CUDA -extern "C" void deviceSearchFilter(PsiPhiArray& psi_phi_data, float* image_times, SearchParameters params, - int num_trajectories, Trajectory* trj_to_search, int num_results, - Trajectory* best_results); +extern "C" void deviceSearchFilter(SearchData &search_data, SearchParameters params, int num_trajectories, + Trajectory *trj_to_search, int num_results, Trajectory *best_results); #endif StackSearch::StackSearch(ImageStack& imstack) : stack(imstack) { @@ -76,8 +75,8 @@ void StackSearch::search(int ang_steps, int vel_steps, float min_ang, float max_ DebugTimer psi_phi_timer = DebugTimer("Creating psi/phi buffers", debug_info); prepare_psi_phi(); - PsiPhiArray psi_phi_data; - fill_psi_phi_array(psi_phi_data, params.encode_num_bytes, psi_images, phi_images, debug_info); + SearchData psi_phi_data; + fill_search_data(psi_phi_data, params.encode_num_bytes, psi_images, phi_images, debug_info); psi_phi_timer.stop(); // Allocate a vector for the results. @@ -98,7 +97,7 @@ void StackSearch::search(int ang_steps, int vel_steps, float min_ang, float max_ // Do the actual search on the GPU. DebugTimer search_timer = DebugTimer("Running search", debug_info); #ifdef HAVE_CUDA - deviceSearchFilter(psi_phi_data, image_times.data(), params, search_list.size(), search_list.data(), + deviceSearchFilter(psi_phi_data, params, search_list.size(), search_list.data(), max_results, results.data()); #else throw std::runtime_error("Non-GPU search is not implemented."); From 881a058f466acd19982b1b5f60bd4e94da87d44a Mon Sep 17 00:00:00 2001 From: Jeremy Kubica <104161096+jeremykubica@users.noreply.github.com> Date: Thu, 1 Feb 2024 11:14:43 -0500 Subject: [PATCH 02/11] File rename --- .../search/pydocs/{psi_phi_array_docs.h => search_data_docs.h} | 0 1 file changed, 0 insertions(+), 0 deletions(-) rename src/kbmod/search/pydocs/{psi_phi_array_docs.h => search_data_docs.h} (100%) diff --git a/src/kbmod/search/pydocs/psi_phi_array_docs.h b/src/kbmod/search/pydocs/search_data_docs.h similarity index 100% rename from src/kbmod/search/pydocs/psi_phi_array_docs.h rename to src/kbmod/search/pydocs/search_data_docs.h From ce99d654d67b6afc66c3237340979918122f4e33 Mon Sep 17 00:00:00 2001 From: Jeremy Kubica <104161096+jeremykubica@users.noreply.github.com> Date: Thu, 1 Feb 2024 11:29:48 -0500 Subject: [PATCH 03/11] Add tests --- src/kbmod/search/search_data_utils.h | 2 +- ...t_psi_phi_array.py => test_search_data.py} | 77 +++++++++++++++++-- 2 files changed, 70 insertions(+), 9 deletions(-) rename tests/{test_psi_phi_array.py => test_search_data.py} (72%) diff --git a/src/kbmod/search/search_data_utils.h b/src/kbmod/search/search_data_utils.h index 50538e99..4d300af6 100644 --- a/src/kbmod/search/search_data_utils.h +++ b/src/kbmod/search/search_data_utils.h @@ -32,7 +32,7 @@ void fill_search_data(SearchData& result_data, int num_bytes, const std::vector< bool debug = false); void fill_search_data_from_image_stack(SearchData& result_data, ImageStack& stack, int num_bytes, - bool debug = false); + bool debug = false); } /* namespace search */ diff --git a/tests/test_psi_phi_array.py b/tests/test_search_data.py similarity index 72% rename from tests/test_psi_phi_array.py rename to tests/test_search_data.py index 8260fffa..6e3e234e 100644 --- a/tests/test_psi_phi_array.py +++ b/tests/test_search_data.py @@ -3,18 +3,23 @@ import numpy as np from kbmod.search import ( + HAS_GPU, KB_NO_DATA, + PSF, + ImageStack, + LayeredImage, PsiPhi, - PsiPhiArray, + SearchData, RawImage, compute_scale_params_from_image_vect, decode_uint_scalar, encode_uint_scalar, - fill_psi_phi_array, + fill_search_data, + fill_search_data_from_image_stack, ) -class test_psi_phi_array(unittest.TestCase): +class test_search_data(unittest.TestCase): def setUp(self): self.num_times = 2 self.width = 4 @@ -31,8 +36,10 @@ def setUp(self): self.phi_1 = RawImage(np.full((self.height, self.width), 0.1, dtype=np.single), obs_time=1.0) self.phi_2 = RawImage(np.full((self.height, self.width), 0.2, dtype=np.single), obs_time=2.0) + self.zeroed_times = [0.0, 1.0] + def test_set_meta_data(self): - arr = PsiPhiArray() + arr = SearchData() self.assertEqual(arr.num_times, 0) self.assertEqual(arr.num_bytes, 4) self.assertEqual(arr.width, 0) @@ -121,10 +128,10 @@ def test_compute_scale_params_from_image_vect(self): self.assertAlmostEqual(result_uint16[1], max_val, delta=1e-5) self.assertAlmostEqual(result_uint16[2], max_val / 65535.0, delta=1e-5) - def test_fill_psi_phi_array(self): + def test_fill_search_data(self): for num_bytes in [2, 4]: - arr = PsiPhiArray() - fill_psi_phi_array(arr, num_bytes, [self.psi_1, self.psi_2], [self.phi_1, self.phi_2], False) + arr = SearchData() + fill_search_data(arr, num_bytes, [self.psi_1, self.psi_2], [self.phi_1, self.phi_2], self.zeroed_times, False) # Check the meta data. self.assertEqual(arr.num_times, self.num_times) @@ -139,9 +146,16 @@ def test_fill_psi_phi_array(self): self.assertEqual(arr.block_size, num_bytes) self.assertEqual(arr.total_array_size, arr.num_entries * arr.block_size) - # Check that we can correctly read the values from the CPU. + # Check that we allocate the arrays self.assertTrue(arr.cpu_array_allocated) + self.assertTrue(arr.cpu_time_array_allocated) + if (HAS_GPU): + self.assertTrue(arr.gpu_array_allocated) + self.assertTrue(arr.gpu_time_array_allocated) + + # Check that we can correctly read the values from the CPU. for time in range(self.num_times): + self.assertAlmostEqual(arr.read_time(time), self.zeroed_times[time]) offset = time * self.width * self.height for row in range(self.height): for col in range(self.width): @@ -152,6 +166,53 @@ def test_fill_psi_phi_array(self): # Check that the arrays are set to NULL after we clear it (memory should be freed too). arr.clear() self.assertFalse(arr.cpu_array_allocated) + self.assertFalse(arr.cpu_time_array_allocated) + if (HAS_GPU): + self.assertFalse(arr.gpu_array_allocated) + self.assertFalse(arr.gpu_time_array_allocated) + + def test_fill_search_data_from_image_stack(self): + # Build a fake image stack. + num_images = 5 + width = 21 + height = 15 + images = [None] * num_images + p = PSF(1.0) + for i in range(num_images): + self.images[i] = kb.LayeredImage( + width, + height, + 2.0, # noise_level + 4.0, # variance + 2.0 * i + 1.0, # time + p, + ) + im_stack = ImageStack(images) + + # Create the SearchData from the ImageStack. + arr = SearchData() + fill_search_data_from_image_stack(arr, im_stack, 4, False) + + # Check the meta data. + self.assertEqual(arr.num_times, num_images) + self.assertEqual(arr.num_bytes, 4) + self.assertEqual(arr.width, width) + self.assertEqual(arr.height, height) + self.assertEqual(arr.pixels_per_image, width * height) + self.assertEqual(arr.num_entries, 2 * arr.pixels_per_image * num_times) + self.assertEqual(arr.block_size, 4) + self.assertEqual(arr.total_array_size, arr.num_entries * arr.block_size) + + # Check that we allocated the arrays. + self.assertTrue(arr.cpu_array_allocated) + self.assertTrue(arr.cpu_time_array_allocated) + if (HAS_GPU): + self.assertTrue(arr.gpu_array_allocated) + self.assertTrue(arr.gpu_time_array_allocated) + + # Since we filled the images with random data, we only test the times. + for time in range(self.num_times): + self.assertAlmostEqual(arr.read_time(time), 2.0 * time) if __name__ == "__main__": From 0b2b94124926e4776757561a1cc77d99fc8870a8 Mon Sep 17 00:00:00 2001 From: Jeremy Kubica <104161096+jeremykubica@users.noreply.github.com> Date: Thu, 1 Feb 2024 12:13:19 -0500 Subject: [PATCH 04/11] Bug fixes and linting --- src/kbmod/search/bindings.cpp | 4 +-- src/kbmod/search/kernels.cu | 12 ++++---- src/kbmod/search/search_data.cpp | 41 ++++++++++++++-------------- src/kbmod/search/search_data_ds.h | 9 +++--- src/kbmod/search/search_data_utils.h | 4 +-- src/kbmod/search/stack_search.cpp | 10 +++---- src/kbmod/search/stack_search.h | 4 +-- tests/test_search_data.py | 24 ++++++++-------- 8 files changed, 56 insertions(+), 52 deletions(-) diff --git a/src/kbmod/search/bindings.cpp b/src/kbmod/search/bindings.cpp index 4175c963..a76eb064 100644 --- a/src/kbmod/search/bindings.cpp +++ b/src/kbmod/search/bindings.cpp @@ -16,7 +16,7 @@ namespace py = pybind11; #include "stack_search.cpp" #include "stamp_creator.cpp" #include "kernel_testing_helpers.cpp" -#include "psi_phi_array.cpp" +#include "search_data.cpp" PYBIND11_MODULE(search, m) { m.attr("KB_NO_DATA") = pybind11::float_(search::NO_DATA); @@ -40,7 +40,7 @@ PYBIND11_MODULE(search, m) { search::pixel_pos_bindings(m); search::image_moments_bindings(m); search::stamp_parameters_bindings(m); - search::psi_phi_array_binding(m); + search::search_data_binding(m); // Functions from raw_image.cpp m.def("create_median_image", &search::create_median_image); m.def("create_summed_image", &search::create_summed_image); diff --git a/src/kbmod/search/kernels.cu b/src/kbmod/search/kernels.cu index 2accdb17..c5036a0a 100644 --- a/src/kbmod/search/kernels.cu +++ b/src/kbmod/search/kernels.cu @@ -42,7 +42,8 @@ extern "C" void device_allocate_search_data_arrays(SearchData *data) { float *device_times_ptr; long unsigned time_bytes = data->get_num_times() * sizeof(float); checkCudaErrors(cudaMalloc((void **)&device_times_ptr, time_bytes)); - checkCudaErrors(cudaMemcpy(device_times_ptr, data->get_cpu_time_array_ptr(), time_bytes, cudaMemcpyHostToDevice)); + checkCudaErrors( + cudaMemcpy(device_times_ptr, data->get_cpu_time_array_ptr(), time_bytes, cudaMemcpyHostToDevice)); data->set_gpu_time_array_ptr(device_times_ptr); } @@ -296,7 +297,7 @@ extern "C" void deviceSearchFilter(SearchData &search_data, SearchParameters par // Copy trajectories to search if (params.debug) { printf("Allocating GPU memory for testing grid with %i elements using %lu bytes.\n", num_trajectories, - sizeof(Trajectory) * num_trajectories); + sizeof(Trajectory) * num_trajectories); } checkCudaErrors(cudaMalloc((void **)&device_tests, sizeof(Trajectory) * num_trajectories)); checkCudaErrors(cudaMemcpy(device_tests, trj_to_search, sizeof(Trajectory) * num_trajectories, @@ -304,7 +305,8 @@ extern "C" void deviceSearchFilter(SearchData &search_data, SearchParameters par // Allocate space for the results. if (params.debug) { - printf("Allocating GPU memory for %i results using %lu bytes.\n", num_results, sizeof(Trajectory) * num_results); + printf("Allocating GPU memory for %i results using %lu bytes.\n", num_results, + sizeof(Trajectory) * num_results); } checkCudaErrors(cudaMalloc((void **)&device_search_results, sizeof(Trajectory) * num_results)); @@ -318,8 +320,8 @@ extern "C" void deviceSearchFilter(SearchData &search_data, SearchParameters par // Launch Search searchFilterImages<<>>(search_data.get_meta_data(), search_data.get_gpu_array_ptr(), - static_cast(search_data.get_gpu_time_array_ptr()), params, - num_trajectories, device_tests, device_search_results); + static_cast(search_data.get_gpu_time_array_ptr()), + params, num_trajectories, device_tests, device_search_results); // Read back results checkCudaErrors(cudaMemcpy(best_results, device_search_results, sizeof(Trajectory) * num_results, diff --git a/src/kbmod/search/search_data.cpp b/src/kbmod/search/search_data.cpp index 4331f779..ae2e1c42 100644 --- a/src/kbmod/search/search_data.cpp +++ b/src/kbmod/search/search_data.cpp @@ -17,9 +17,7 @@ extern "C" void device_free_search_data_arrays(SearchData* data); SearchData::SearchData() {} -SearchData::~SearchData() { - clear(); -} +SearchData::~SearchData() { clear(); } void SearchData::clear() { // Free all used memory on CPU and GPU. @@ -29,7 +27,7 @@ void SearchData::clear() { } if (cpu_time_array != nullptr) { free(cpu_time_array); - cpu_array_ptr = nullptr; + cpu_time_array = nullptr; } #ifdef HAVE_CUDA if ((gpu_array_ptr != nullptr) || (gpu_time_array != nullptr)) { @@ -140,13 +138,12 @@ PsiPhi SearchData::read_psi_phi(int time, int row, int col) { float SearchData::read_time(int time_index) { if (cpu_time_array == nullptr) throw std::runtime_error("Read from unallocated times array."); - if ((time_index < 0 )|| (time_index >= meta_data.num_times)) { - throw std::runtime_error("Out of bounds read for time step=%i", time_index); + if ((time_index < 0) || (time_index >= meta_data.num_times)) { + throw std::runtime_error("Out of bounds read for time step."); } return cpu_time_array[time_index]; } - // ------------------------------------------- // --- Implementation of utility functions --- // ------------------------------------------- @@ -179,7 +176,7 @@ std::array compute_scale_params_from_image_vect(const std::vector void set_encode_cpu_search_data(SearchData& data, const std::vector& psi_imgs, - const std::vector& phi_imgs, bool debug) { + const std::vector& phi_imgs, bool debug) { if (data.get_cpu_array_ptr() != nullptr) { throw std::runtime_error("CPU PsiPhi already allocated."); } @@ -223,7 +220,7 @@ void set_encode_cpu_search_data(SearchData& data, const std::vector& p } void set_float_cpu_search_data(SearchData& data, const std::vector& psi_imgs, - const std::vector& phi_imgs, bool debug) { + const std::vector& phi_imgs, bool debug) { if (data.get_cpu_array_ptr() != nullptr) { throw std::runtime_error("CPU PsiPhi already allocated."); } @@ -249,8 +246,8 @@ void set_float_cpu_search_data(SearchData& data, const std::vector& ps } void fill_search_data(SearchData& result_data, int num_bytes, const std::vector& psi_imgs, - const std::vector& phi_imgs, const std::vector zeroed_times, - bool debug) { + const std::vector& phi_imgs, const std::vector zeroed_times, + bool debug) { if (result_data.get_cpu_array_ptr() != nullptr) { return; } @@ -259,7 +256,8 @@ void fill_search_data(SearchData& result_data, int num_bytes, const std::vector< int num_times = psi_imgs.size(); if (num_times <= 0) throw std::runtime_error("Trying to fill PsiPhi from empty vectors."); if (num_times != phi_imgs.size()) throw std::runtime_error("Size mismatch between psi and phi."); - if (num_times != zeroed_times.size()) throw std::runtime_error("Size mismatch between psi and zeroed times."); + if (num_times != zeroed_times.size()) + throw std::runtime_error("Size mismatch between psi and zeroed times."); int width = phi_imgs[0].get_width(); int height = phi_imgs[0].get_height(); @@ -298,14 +296,14 @@ void fill_search_data(SearchData& result_data, int num_bytes, const std::vector< // Copy the time array. const long unsigned times_bytes = result_data.get_num_times() * sizeof(float); - if (debug) printf("Allocating %lu bytes on the CPU for times.\n"); + if (debug) printf("Allocating %lu bytes on the CPU for times.\n", times_bytes); - float* times_array = (float*)malloc(data.get_total_array_size()); + float* times_array = (float*)malloc(times_bytes); if (times_array == nullptr) throw std::runtime_error("Unable to allocate space for CPU times."); for (int i = 0; i < result_data.get_num_times(); ++i) { times_array[i] = zeroed_times[i]; } - data.set_cpu_time_array_ptr((void*)times_array); + result_data.set_cpu_time_array_ptr(times_array); #ifdef HAVE_CUDA // Create a copy of the encoded data in GPU memory. @@ -314,7 +312,7 @@ void fill_search_data(SearchData& result_data, int num_bytes, const std::vector< result_data.get_total_array_size()); printf("Allocating GPU memory for times array using %lu bytes.\n", times_bytes); } - + device_allocate_search_data_arrays(&result_data); if (result_data.get_gpu_array_ptr() == nullptr) { throw std::runtime_error("Unable to allocate GPU PsiPhi array."); @@ -325,7 +323,8 @@ void fill_search_data(SearchData& result_data, int num_bytes, const std::vector< #endif } -void fill_search_data_from_image_stack(SearchData& result_data, ImageStack& stack, int num_bytes, bool debug) { +void fill_search_data_from_image_stack(SearchData& result_data, ImageStack& stack, int num_bytes, + bool debug) { // Compute Phi and Psi from convolved images while leaving masked pixels alone // Reinsert 0s for NO_DATA? std::vector psi_images; @@ -333,8 +332,8 @@ void fill_search_data_from_image_stack(SearchData& result_data, ImageStack& stac const int num_images = stack.img_count(); if (debug) { unsigned long num_bytes = 2 * stack.get_height() * stack.get_width() * num_images * sizeof(float); - printf("Building %i temporary %i by %i images (psi and phi), requiring %lu bytes", - (num_images * 2), stack.get_width(), stack.get_height(), num_bytes); + printf("Building %i temporary %i by %i images (psi and phi), requiring %lu bytes", (num_images * 2), + stack.get_width(), stack.get_height(), num_bytes); } // Build the psi and phi images first. @@ -347,7 +346,7 @@ void fill_search_data_from_image_stack(SearchData& result_data, ImageStack& stac // Convert these into an array form. Needs the full psi and phi computed first so the // encoding can compute the bounds of each array. std::vector zeroed_times = stack.build_zeroed_times(); - fill_search_data(psi_phi_data, num_bytes, psi_images, phi_images, zeroed_times, debug); + fill_search_data(result_data, num_bytes, psi_images, phi_images, zeroed_times, debug); } // ------------------------------------------- @@ -396,7 +395,7 @@ static void search_data_binding(py::module& m) { pydocs::DOC_SearchData_get_gpu_time_array_allocated) .def("set_meta_data", &ppa::set_meta_data, pydocs::DOC_SearchData_set_meta_data) .def("clear", &ppa::clear, pydocs::DOC_SearchData_clear) - .def("read_psi_phi", &ppa::read_psi_phi, pydocs::DOC_SearchData_read_psi_phi); + .def("read_psi_phi", &ppa::read_psi_phi, pydocs::DOC_SearchData_read_psi_phi) .def("read_time", &ppa::read_time, pydocs::DOC_SearchData_read_time); m.def("compute_scale_params_from_image_vect", &search::compute_scale_params_from_image_vect); m.def("decode_uint_scalar", &search::decode_uint_scalar); diff --git a/src/kbmod/search/search_data_ds.h b/src/kbmod/search/search_data_ds.h index 900c3cf5..7a64842c 100644 --- a/src/kbmod/search/search_data_ds.h +++ b/src/kbmod/search/search_data_ds.h @@ -8,10 +8,11 @@ * only what they need. * * The data structure allocates memory on both the CPU and GPU for the - * arraysand maintains ownership of the pointers until clear() is called + * arrays and maintains ownership of the pointers until clear() is called * the object's destructor is called. This allows the object to be passed * repeatedly to the on-device search without reallocating and copying the - * memory on the GPU. + * memory on the GPU. All arrays are stored as pointers (instead of vectors) + * for compatibility with CUDA. * * Created on: Dec 5, 2023 */ @@ -101,7 +102,7 @@ class SearchData { // Primary getter functions for interaction (read the data). PsiPhi read_psi_phi(int time_index, int row, int col); - float read_time_value(int time_index); + float read_time(int time_index); // Setters for the utility functions to allocate the data. void set_meta_data(int new_num_bytes, int new_num_times, int new_height, int new_width); @@ -123,7 +124,7 @@ class SearchData { SearchDataMeta meta_data; // Pointers to the arrays - void* cpu_array_ptr = nullptr; + void* cpu_array_ptr = nullptr; void* gpu_array_ptr = nullptr; float* cpu_time_array = nullptr; float* gpu_time_array = nullptr; diff --git a/src/kbmod/search/search_data_utils.h b/src/kbmod/search/search_data_utils.h index 4d300af6..683895ce 100644 --- a/src/kbmod/search/search_data_utils.h +++ b/src/kbmod/search/search_data_utils.h @@ -28,8 +28,8 @@ namespace search { std::array compute_scale_params_from_image_vect(const std::vector& imgs, int num_bytes); void fill_search_data(SearchData& result_data, int num_bytes, const std::vector& psi_imgs, - const std::vector& phi_imgs, const std::vector zeroed_times, - bool debug = false); + const std::vector& phi_imgs, const std::vector zeroed_times, + bool debug = false); void fill_search_data_from_image_stack(SearchData& result_data, ImageStack& stack, int num_bytes, bool debug = false); diff --git a/src/kbmod/search/stack_search.cpp b/src/kbmod/search/stack_search.cpp index 60ff1611..64c6741b 100644 --- a/src/kbmod/search/stack_search.cpp +++ b/src/kbmod/search/stack_search.cpp @@ -2,8 +2,8 @@ namespace search { #ifdef HAVE_CUDA -extern "C" void deviceSearchFilter(SearchData &search_data, SearchParameters params, int num_trajectories, - Trajectory *trj_to_search, int num_results, Trajectory *best_results); +extern "C" void deviceSearchFilter(SearchData& search_data, SearchParameters params, int num_trajectories, + Trajectory* trj_to_search, int num_results, Trajectory* best_results); #endif StackSearch::StackSearch(ImageStack& imstack) : stack(imstack) { @@ -76,7 +76,7 @@ void StackSearch::search(int ang_steps, int vel_steps, float min_ang, float max_ DebugTimer psi_phi_timer = DebugTimer("Creating psi/phi buffers", debug_info); prepare_psi_phi(); SearchData psi_phi_data; - fill_search_data(psi_phi_data, params.encode_num_bytes, psi_images, phi_images, debug_info); + fill_search_data(psi_phi_data, params.encode_num_bytes, psi_images, phi_images, image_times, debug_info); psi_phi_timer.stop(); // Allocate a vector for the results. @@ -97,8 +97,8 @@ void StackSearch::search(int ang_steps, int vel_steps, float min_ang, float max_ // Do the actual search on the GPU. DebugTimer search_timer = DebugTimer("Running search", debug_info); #ifdef HAVE_CUDA - deviceSearchFilter(psi_phi_data, params, search_list.size(), search_list.data(), - max_results, results.data()); + deviceSearchFilter(psi_phi_data, params, search_list.size(), search_list.data(), max_results, + results.data()); #else throw std::runtime_error("Non-GPU search is not implemented."); #endif diff --git a/src/kbmod/search/stack_search.h b/src/kbmod/search/stack_search.h index 597f7885..5b5af451 100644 --- a/src/kbmod/search/stack_search.h +++ b/src/kbmod/search/stack_search.h @@ -16,8 +16,8 @@ #include "geom.h" #include "image_stack.h" #include "psf.h" -#include "psi_phi_array_ds.h" -#include "psi_phi_array_utils.h" +#include "search_data_ds.h" +#include "search_data_utils.h" #include "pydocs/stack_search_docs.h" #include "stamp_creator.h" diff --git a/tests/test_search_data.py b/tests/test_search_data.py index 6e3e234e..80b84736 100644 --- a/tests/test_search_data.py +++ b/tests/test_search_data.py @@ -131,7 +131,9 @@ def test_compute_scale_params_from_image_vect(self): def test_fill_search_data(self): for num_bytes in [2, 4]: arr = SearchData() - fill_search_data(arr, num_bytes, [self.psi_1, self.psi_2], [self.phi_1, self.phi_2], self.zeroed_times, False) + fill_search_data( + arr, num_bytes, [self.psi_1, self.psi_2], [self.phi_1, self.phi_2], self.zeroed_times, False + ) # Check the meta data. self.assertEqual(arr.num_times, self.num_times) @@ -149,7 +151,7 @@ def test_fill_search_data(self): # Check that we allocate the arrays self.assertTrue(arr.cpu_array_allocated) self.assertTrue(arr.cpu_time_array_allocated) - if (HAS_GPU): + if HAS_GPU: self.assertTrue(arr.gpu_array_allocated) self.assertTrue(arr.gpu_time_array_allocated) @@ -167,19 +169,19 @@ def test_fill_search_data(self): arr.clear() self.assertFalse(arr.cpu_array_allocated) self.assertFalse(arr.cpu_time_array_allocated) - if (HAS_GPU): + if HAS_GPU: self.assertFalse(arr.gpu_array_allocated) self.assertFalse(arr.gpu_time_array_allocated) def test_fill_search_data_from_image_stack(self): # Build a fake image stack. - num_images = 5 + num_times = 5 width = 21 height = 15 - images = [None] * num_images + images = [None] * num_times p = PSF(1.0) - for i in range(num_images): - self.images[i] = kb.LayeredImage( + for i in range(num_times): + images[i] = LayeredImage( width, height, 2.0, # noise_level @@ -192,9 +194,9 @@ def test_fill_search_data_from_image_stack(self): # Create the SearchData from the ImageStack. arr = SearchData() fill_search_data_from_image_stack(arr, im_stack, 4, False) - + # Check the meta data. - self.assertEqual(arr.num_times, num_images) + self.assertEqual(arr.num_times, num_times) self.assertEqual(arr.num_bytes, 4) self.assertEqual(arr.width, width) self.assertEqual(arr.height, height) @@ -206,12 +208,12 @@ def test_fill_search_data_from_image_stack(self): # Check that we allocated the arrays. self.assertTrue(arr.cpu_array_allocated) self.assertTrue(arr.cpu_time_array_allocated) - if (HAS_GPU): + if HAS_GPU: self.assertTrue(arr.gpu_array_allocated) self.assertTrue(arr.gpu_time_array_allocated) # Since we filled the images with random data, we only test the times. - for time in range(self.num_times): + for time in range(num_times): self.assertAlmostEqual(arr.read_time(time), 2.0 * time) From 0839114ec797a8869f94cbd5e5821a9b1e3d9eb2 Mon Sep 17 00:00:00 2001 From: Jeremy Kubica <104161096+jeremykubica@users.noreply.github.com> Date: Fri, 2 Feb 2024 07:58:44 -0500 Subject: [PATCH 05/11] Fix bad merge --- src/kbmod/search/kernels.cu | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/src/kbmod/search/kernels.cu b/src/kbmod/search/kernels.cu index a70e41bc..21bf19a9 100644 --- a/src/kbmod/search/kernels.cu +++ b/src/kbmod/search/kernels.cu @@ -337,8 +337,8 @@ extern "C" void deviceSearchFilter(SearchData &search_data, SearchParameters par // Launch Search searchFilterImages<<>>(search_data.get_meta_data(), search_data.get_gpu_array_ptr(), - device_img_times, params, num_trajectories, device_tests, - device_search_results); + static_cast(search_data.get_gpu_time_array_ptr()), + params, num_trajectories, device_tests, device_search_results); cudaDeviceSynchronize(); // Read back results From 43e82a83d5e811b3f7978ba6b24c5863bd6af0ad Mon Sep 17 00:00:00 2001 From: Jeremy Kubica <104161096+jeremykubica@users.noreply.github.com> Date: Fri, 2 Feb 2024 16:20:57 -0500 Subject: [PATCH 06/11] Revert the name change --- src/kbmod/search/bindings.cpp | 4 +- src/kbmod/search/kernels.cu | 31 ++++----- src/kbmod/search/layered_image.cpp | 8 +-- .../{search_data.cpp => psi_phi_array.cpp} | 0 .../{search_data_ds.h => psi_phi_array_ds.h} | 24 +++---- ...rch_data_utils.h => psi_phi_array_utils.h} | 20 +++--- src/kbmod/search/pydocs/search_data_docs.h | 66 +++++++++---------- src/kbmod/search/stack_search.cpp | 7 +- src/kbmod/search/stack_search.h | 4 +- ...t_search_data.py => test_psi_phi_array.py} | 10 +-- 10 files changed, 88 insertions(+), 86 deletions(-) rename src/kbmod/search/{search_data.cpp => psi_phi_array.cpp} (100%) rename src/kbmod/search/{search_data_ds.h => psi_phi_array_ds.h} (91%) rename src/kbmod/search/{search_data_utils.h => psi_phi_array_utils.h} (52%) rename tests/{test_search_data.py => test_psi_phi_array.py} (98%) diff --git a/src/kbmod/search/bindings.cpp b/src/kbmod/search/bindings.cpp index c39343c2..225bf477 100644 --- a/src/kbmod/search/bindings.cpp +++ b/src/kbmod/search/bindings.cpp @@ -16,7 +16,7 @@ namespace py = pybind11; #include "stack_search.cpp" #include "stamp_creator.cpp" #include "kernel_testing_helpers.cpp" -#include "search_data.cpp" +#include "psi_phi_array.cpp" PYBIND11_MODULE(search, m) { m.attr("KB_NO_DATA") = pybind11::float_(search::NO_DATA); @@ -39,7 +39,7 @@ PYBIND11_MODULE(search, m) { search::trajectory_bindings(m); search::image_moments_bindings(m); search::stamp_parameters_bindings(m); - search::search_data_binding(m); + search::psi_phi_array_binding(m); // Functions from raw_image.cpp m.def("create_median_image", &search::create_median_image); m.def("create_summed_image", &search::create_summed_image); diff --git a/src/kbmod/search/kernels.cu b/src/kbmod/search/kernels.cu index 21bf19a9..c400c438 100644 --- a/src/kbmod/search/kernels.cu +++ b/src/kbmod/search/kernels.cu @@ -19,11 +19,11 @@ #include "common.h" #include "cuda_errors.h" -#include "search_data_ds.h" +#include "psi_phi_array_ds.h" namespace search { -extern "C" void device_allocate_search_data_arrays(SearchData *data) { +extern "C" void device_allocate_psi_phi_array_arrays(PsiPhiArray *data) { if (!data->cpu_array_allocated() || !data->cpu_time_array_allocated()) { throw std::runtime_error("CPU data is not allocated."); } @@ -47,7 +47,7 @@ extern "C" void device_allocate_search_data_arrays(SearchData *data) { data->set_gpu_time_array_ptr(device_times_ptr); } -extern "C" void device_free_search_data_arrays(SearchData *data) { +extern "C" void device_free_psi_phi_array_arrays(PsiPhiArray *data) { if (data->gpu_array_allocated()) { checkCudaErrors(cudaFree(data->get_gpu_array_ptr())); data->set_gpu_array_ptr(nullptr); @@ -58,8 +58,8 @@ extern "C" void device_free_search_data_arrays(SearchData *data) { } } -__host__ __device__ PsiPhi read_encoded_psi_phi(SearchDataMeta ¶ms, void *psi_phi_vect, int time, - int row, int col) { +__host__ __device__ PsiPhi read_encoded_psi_phi(PsiPhiArrayMeta ¶ms, void *psi_phi_vect, int time, + int row, int col) { // Bounds checking. if ((row < 0) || (col < 0) || (row >= params.height) || (col >= params.width)) { return {NO_DATA, NO_DATA}; @@ -141,11 +141,12 @@ extern "C" __device__ __host__ void SigmaGFilteredIndicesCU(float *values, int n /* * Evaluate the likelihood score (as computed with from the psi and phi values) for a single - * given candidate trajectory. Modifies the trajectory in place to update the number of + * given candidate trajectory. Modifies the trajectory in place to update the number of * observations, likelihood, and flux. */ -extern "C" __device__ __host__ void evaluateTrajectory(SearchDataMeta psi_phi_meta, void *psi_phi_vect, float *image_times, - SearchParameters params, Trajectory *candidate) { +extern "C" __device__ __host__ void evaluateTrajectory(PsiPhiArrayMeta psi_phi_meta, void *psi_phi_vect, + float *image_times, SearchParameters params, + Trajectory *candidate) { // Data structures used for filtering. We fill in only what we need. float psi_array[MAX_NUM_IMAGES]; float phi_array[MAX_NUM_IMAGES]; @@ -222,7 +223,7 @@ extern "C" __device__ __host__ void evaluateTrajectory(SearchDataMeta psi_phi_me * * Creates a local copy of psi_phi_meta and params in local memory space. */ -__global__ void searchFilterImages(SearchDataMeta psi_phi_meta, void *psi_phi_vect, float *image_times, +__global__ void searchFilterImages(PsiPhiArrayMeta psi_phi_meta, void *psi_phi_vect, float *image_times, SearchParameters params, int num_trajectories, Trajectory *trajectories, Trajectory *results) { // Get the x and y coordinates within the search space. @@ -291,23 +292,23 @@ __global__ void searchFilterImages(SearchDataMeta psi_phi_meta, void *psi_phi_ve } } -extern "C" void deviceSearchFilter(SearchData &search_data, SearchParameters params, int num_trajectories, +extern "C" void deviceSearchFilter(PsiPhiArray &psi_phi_array, SearchParameters params, int num_trajectories, Trajectory *trj_to_search, int num_results, Trajectory *best_results) { // Allocate Device memory Trajectory *device_tests; Trajectory *device_search_results; // Check the hard coded maximum number of images against the num_images. - int num_images = search_data.get_num_times(); + int num_images = psi_phi_array.get_num_times(); if (num_images > MAX_NUM_IMAGES) { throw std::runtime_error("Number of images exceeds GPU maximum."); } // Check that the device vectors have already been allocated. - if (search_data.gpu_array_allocated() == false) { + if (psi_phi_array.gpu_array_allocated() == false) { throw std::runtime_error("PsiPhi data has not been created."); } - if (search_data.gpu_time_array_allocated() == false) { + if (psi_phi_array.gpu_time_array_allocated() == false) { throw std::runtime_error("GPU time data has not been created."); } @@ -336,8 +337,8 @@ extern "C" void deviceSearchFilter(SearchData &search_data, SearchParameters par dim3 threads(THREAD_DIM_X, THREAD_DIM_Y); // Launch Search - searchFilterImages<<>>(search_data.get_meta_data(), search_data.get_gpu_array_ptr(), - static_cast(search_data.get_gpu_time_array_ptr()), + searchFilterImages<<>>(psi_phi_array.get_meta_data(), psi_phi_array.get_gpu_array_ptr(), + static_cast(psi_phi_array.get_gpu_time_array_ptr()), params, num_trajectories, device_tests, device_search_results); cudaDeviceSynchronize(); diff --git a/src/kbmod/search/layered_image.cpp b/src/kbmod/search/layered_image.cpp index 903ec82a..ba7efbb0 100644 --- a/src/kbmod/search/layered_image.cpp +++ b/src/kbmod/search/layered_image.cpp @@ -40,12 +40,12 @@ LayeredImage::LayeredImage(const RawImage& sci, const RawImage& var, const RawIm variance = var; } -LayeredImage::LayeredImage(unsigned w, unsigned h, float noise_stdev, float pixel_variance, - double time, const PSF& psf) +LayeredImage::LayeredImage(unsigned w, unsigned h, float noise_stdev, float pixel_variance, double time, + const PSF& psf) : LayeredImage(w, h, noise_stdev, pixel_variance, time, psf, -1) {} -LayeredImage::LayeredImage(unsigned w, unsigned h, float noise_stdev, float pixel_variance, - double time, const PSF& psf, int seed) +LayeredImage::LayeredImage(unsigned w, unsigned h, float noise_stdev, float pixel_variance, double time, + const PSF& psf, int seed) : psf(psf), width(w), height(h) { std::random_device r; std::default_random_engine generator(r()); diff --git a/src/kbmod/search/search_data.cpp b/src/kbmod/search/psi_phi_array.cpp similarity index 100% rename from src/kbmod/search/search_data.cpp rename to src/kbmod/search/psi_phi_array.cpp diff --git a/src/kbmod/search/search_data_ds.h b/src/kbmod/search/psi_phi_array_ds.h similarity index 91% rename from src/kbmod/search/search_data_ds.h rename to src/kbmod/search/psi_phi_array_ds.h index 7a64842c..477b1704 100644 --- a/src/kbmod/search/search_data_ds.h +++ b/src/kbmod/search/psi_phi_array_ds.h @@ -1,5 +1,5 @@ /* - * search_data_ds.h + * psi_phi_array_ds.h * * The data structure for the raw data needed for the search algorith, * including the psi/phi values and the zeroed times. The the data @@ -17,8 +17,8 @@ * Created on: Dec 5, 2023 */ -#ifndef SEARCH_DATA_DS_ -#define SEARCH_DATA_DS_ +#ifndef PSI_PHI_ARRAY_DS_ +#define PSI_PHI_ARRAY_DS_ #include #include @@ -44,8 +44,8 @@ inline float decode_uint_scalar(float value, float min_val, float scale) { return (value == 0.0) ? NO_DATA : (value - 1.0) * scale + min_val; } -// The struct of meta data for the SearchData. -struct SearchDataMeta { +// The struct of meta data for the PsiPhiArray. +struct PsiPhiArrayMeta { int num_times = 0; int width = 0; int height = 0; @@ -66,17 +66,17 @@ struct SearchDataMeta { float phi_scale = 1.0; }; -/* SearchData is a class to hold the psi and phi arrays for the CPU and GPU as well as +/* PsiPhiArray is a class to hold the psi and phi arrays for the CPU and GPU as well as the meta data and functions to do encoding and decoding on CPU. */ -class SearchData { +class PsiPhiArray { public: - explicit SearchData(); - virtual ~SearchData(); + explicit PsiPhiArray(); + virtual ~PsiPhiArray(); void clear(); - inline SearchDataMeta& get_meta_data() { return meta_data; } + inline PsiPhiArrayMeta& get_meta_data() { return meta_data; } // --- Getter functions (for Python interface) ---------------- inline int get_num_bytes() { return meta_data.num_bytes; } @@ -121,7 +121,7 @@ class SearchData { inline void set_gpu_time_array_ptr(float* new_ptr) { gpu_time_array = new_ptr; } private: - SearchDataMeta meta_data; + PsiPhiArrayMeta meta_data; // Pointers to the arrays void* cpu_array_ptr = nullptr; @@ -132,4 +132,4 @@ class SearchData { } /* namespace search */ -#endif /* SEARCH_DATA_DS_ */ +#endif /* PSI_PHI_ARRAY_DS_ */ diff --git a/src/kbmod/search/search_data_utils.h b/src/kbmod/search/psi_phi_array_utils.h similarity index 52% rename from src/kbmod/search/search_data_utils.h rename to src/kbmod/search/psi_phi_array_utils.h index 683895ce..3573df26 100644 --- a/src/kbmod/search/search_data_utils.h +++ b/src/kbmod/search/psi_phi_array_utils.h @@ -1,5 +1,5 @@ /* - * search_data_utils.h + * psi_phi_array_utils.h * * The utility functions for the psi/phi array. Broken out from the header * data structure so that it can use packages that won't be imported into the @@ -8,8 +8,8 @@ * Created on: Dec 8, 2023 */ -#ifndef SEARCH_DATA_UTILS_ -#define SEARCH_DATA_UTILS_ +#ifndef PSI_PHI_ARRAY_UTILS_ +#define PSI_PHI_ARRAY_UTILS_ #include #include @@ -19,7 +19,7 @@ #include "common.h" #include "image_stack.h" #include "layered_image.h" -#include "search_data_ds.h" +#include "psi_phi_array_ds.h" #include "raw_image.h" namespace search { @@ -27,13 +27,13 @@ namespace search { // Compute the min, max, and scale parameter from the a vector of image data. std::array compute_scale_params_from_image_vect(const std::vector& imgs, int num_bytes); -void fill_search_data(SearchData& result_data, int num_bytes, const std::vector& psi_imgs, - const std::vector& phi_imgs, const std::vector zeroed_times, - bool debug = false); +void fill_psi_phi_array(PsiPhiArray& result_data, int num_bytes, const std::vector& psi_imgs, + const std::vector& phi_imgs, const std::vector zeroed_times, + bool debug = false); -void fill_search_data_from_image_stack(SearchData& result_data, ImageStack& stack, int num_bytes, - bool debug = false); +void fill_psi_phi_array_from_image_stack(PsiPhiArray& result_data, ImageStack& stack, int num_bytes, + bool debug = false); } /* namespace search */ -#endif /* SEARCH_DATA_UTILS_ */ +#endif /* PSI_PHI_ARRAY_UTILS_ */ diff --git a/src/kbmod/search/pydocs/search_data_docs.h b/src/kbmod/search/pydocs/search_data_docs.h index 4f668e3e..05abe305 100644 --- a/src/kbmod/search/pydocs/search_data_docs.h +++ b/src/kbmod/search/pydocs/search_data_docs.h @@ -1,5 +1,5 @@ -#ifndef SEARCH_DATA_DOCS -#define SEARCH_DATA_DOCS +#ifndef PSI_PHI_ARRAY_DOCS +#define PSI_PHI_ARRAY_DOCS namespace pydocs { @@ -14,88 +14,88 @@ static const auto DOC_PsiPhi = R"doc( The phi value at a pixel. )doc"; -static const auto DOC_SearchData = R"doc( +static const auto DOC_PsiPhiArray = R"doc( An encoded array of Psi and Phi values along with their meta data. )doc"; -static const auto DOC_SearchData_get_num_bytes = R"doc( +static const auto DOC_PsiPhiArray_get_num_bytes = R"doc( The target number of bytes to use for encoding the data (1 for uint8, 2 for uint16, or 4 for float32). Might differ from actual number of bytes (block_size). )doc"; -static const auto DOC_SearchData_get_num_times = R"doc( +static const auto DOC_PsiPhiArray_get_num_times = R"doc( The number of times. )doc"; -static const auto DOC_SearchData_get_width = R"doc( +static const auto DOC_PsiPhiArray_get_width = R"doc( The image width. )doc"; -static const auto DOC_SearchData_get_height = R"doc( +static const auto DOC_PsiPhiArray_get_height = R"doc( The image height. )doc"; -static const auto DOC_SearchData_get_pixels_per_image = R"doc( +static const auto DOC_PsiPhiArray_get_pixels_per_image = R"doc( The number of pixels per each image. )doc"; -static const auto DOC_SearchData_get_num_entries = R"doc( +static const auto DOC_PsiPhiArray_get_num_entries = R"doc( The number of array entries. )doc"; -static const auto DOC_SearchData_get_total_array_size = R"doc( +static const auto DOC_PsiPhiArray_get_total_array_size = R"doc( The size of the array in bytes. )doc"; -static const auto DOC_SearchData_get_block_size = R"doc( +static const auto DOC_PsiPhiArray_get_block_size = R"doc( The size of a single entry in bytes. )doc"; -static const auto DOC_SearchData_get_psi_min_val = R"doc( +static const auto DOC_PsiPhiArray_get_psi_min_val = R"doc( The minimum value of psi used in the scaling computations. )doc"; -static const auto DOC_SearchData_get_psi_max_val = R"doc( +static const auto DOC_PsiPhiArray_get_psi_max_val = R"doc( The maximum value of psi used in the scaling computations. )doc"; -static const auto DOC_SearchData_get_psi_scale = R"doc( +static const auto DOC_PsiPhiArray_get_psi_scale = R"doc( The scaling parameter for psi. )doc"; -static const auto DOC_SearchData_get_phi_min_val = R"doc( +static const auto DOC_PsiPhiArray_get_phi_min_val = R"doc( The minimum value of phi used in the scaling computations. )doc"; -static const auto DOC_SearchData_get_phi_max_val = R"doc( +static const auto DOC_PsiPhiArray_get_phi_max_val = R"doc( The maximum value of phi used in the scaling computations. )doc"; -static const auto DOC_SearchData_get_phi_scale = R"doc( +static const auto DOC_PsiPhiArray_get_phi_scale = R"doc( The scaling parameter for phi. )doc"; -static const auto DOC_SearchData_get_cpu_array_allocated = R"doc( +static const auto DOC_PsiPhiArray_get_cpu_array_allocated = R"doc( A Boolean indicating whether the cpu data (psi/phi) array exists. )doc"; -static const auto DOC_SearchData_get_gpu_array_allocated = R"doc( +static const auto DOC_PsiPhiArray_get_gpu_array_allocated = R"doc( A Boolean indicating whether the gpu data (psi/phi) array exists. )doc"; -static const auto DOC_SearchData_get_cpu_time_array_allocated = R"doc( +static const auto DOC_PsiPhiArray_get_cpu_time_array_allocated = R"doc( A Boolean indicating whether the cpu time array exists. )doc"; -static const auto DOC_SearchData_get_gpu_time_array_allocated = R"doc( +static const auto DOC_PsiPhiArray_get_gpu_time_array_allocated = R"doc( A Boolean indicating whether the gpu time array exists. )doc"; -static const auto DOC_SearchData_clear = R"doc( +static const auto DOC_PsiPhiArray_clear = R"doc( Clear all data and free the arrays. )doc"; -static const auto DOC_SearchData_read_psi_phi = R"doc( +static const auto DOC_PsiPhiArray_read_psi_phi = R"doc( Read a PsiPhi value from the CPU array. Parameters @@ -113,7 +113,7 @@ static const auto DOC_SearchData_read_psi_phi = R"doc( The pixel values. )doc"; -static const auto DOC_SearchData_read_time = R"doc( +static const auto DOC_PsiPhiArray_read_time = R"doc( Read a zeroed time value from the CPU array. Parameters @@ -127,9 +127,9 @@ static const auto DOC_SearchData_read_time = R"doc( The time. )doc"; -static const auto DOC_SearchData_set_meta_data = R"doc( +static const auto DOC_PsiPhiArray_set_meta_data = R"doc( Set the meta data for the array. Automatically called by - fill_search_data(). + fill_psi_phi_array(). Parameters ---------- @@ -143,12 +143,12 @@ static const auto DOC_SearchData_set_meta_data = R"doc( The width of each image in pixels. )doc"; -static const auto DOC_SearchData_fill_search_data = R"doc( - Fill the SearchData from Psi and Phi images. +static const auto DOC_PsiPhiArray_fill_psi_phi_array = R"doc( + Fill the PsiPhiArray from Psi and Phi images. Parameters ---------- - result_data : `SearchData` + result_data : `PsiPhiArray` The location to store the data. num_bytes : `int` The type of encoding to use (1, 2, or 4). @@ -160,12 +160,12 @@ static const auto DOC_SearchData_fill_search_data = R"doc( A list of floating point times starting at zero. )doc"; -static const auto DOC_SearchData_fill_search_data_from_image_stack = R"doc( - Fill the SearchData an ImageStack. +static const auto DOC_PsiPhiArray_fill_psi_phi_array_from_image_stack = R"doc( + Fill the PsiPhiArray an ImageStack. Parameters ---------- - result_data : `SearchData` + result_data : `PsiPhiArray` The location to store the data. num_bytes : `int` The type of encoding to use (1, 2, or 4). @@ -175,4 +175,4 @@ static const auto DOC_SearchData_fill_search_data_from_image_stack = R"doc( } // namespace pydocs -#endif /* SEARCH_DATA_DOCS */ +#endif /* PSI_PHI_ARRAY_DOCS */ diff --git a/src/kbmod/search/stack_search.cpp b/src/kbmod/search/stack_search.cpp index 401177d6..c945b450 100644 --- a/src/kbmod/search/stack_search.cpp +++ b/src/kbmod/search/stack_search.cpp @@ -2,7 +2,7 @@ namespace search { #ifdef HAVE_CUDA -extern "C" void deviceSearchFilter(SearchData& search_data, SearchParameters params, int num_trajectories, +extern "C" void deviceSearchFilter(PsiPhiArray& psi_phi_array, SearchParameters params, int num_trajectories, Trajectory* trj_to_search, int num_results, Trajectory* best_results); #endif @@ -75,8 +75,9 @@ void StackSearch::search(int ang_steps, int vel_steps, float min_ang, float max_ DebugTimer psi_phi_timer = DebugTimer("Creating psi/phi buffers", debug_info); prepare_psi_phi(); - SearchData psi_phi_data; - fill_search_data(psi_phi_data, params.encode_num_bytes, psi_images, phi_images, image_times, debug_info); + PsiPhiArray psi_phi_data; + fill_psi_phi_array(psi_phi_data, params.encode_num_bytes, psi_images, phi_images, image_times, + debug_info); psi_phi_timer.stop(); // Allocate a vector for the results. diff --git a/src/kbmod/search/stack_search.h b/src/kbmod/search/stack_search.h index ca7ebdd9..fbb9568a 100644 --- a/src/kbmod/search/stack_search.h +++ b/src/kbmod/search/stack_search.h @@ -16,8 +16,8 @@ #include "geom.h" #include "image_stack.h" #include "psf.h" -#include "search_data_ds.h" -#include "search_data_utils.h" +#include "psi_phi_array_ds.h" +#include "psi_phi_array_utils.h" #include "pydocs/stack_search_docs.h" #include "stamp_creator.h" diff --git a/tests/test_search_data.py b/tests/test_psi_phi_array.py similarity index 98% rename from tests/test_search_data.py rename to tests/test_psi_phi_array.py index 80b84736..0e805dd7 100644 --- a/tests/test_search_data.py +++ b/tests/test_psi_phi_array.py @@ -9,7 +9,7 @@ ImageStack, LayeredImage, PsiPhi, - SearchData, + PsiPhiArray, RawImage, compute_scale_params_from_image_vect, decode_uint_scalar, @@ -39,7 +39,7 @@ def setUp(self): self.zeroed_times = [0.0, 1.0] def test_set_meta_data(self): - arr = SearchData() + arr = PsiPhiArray() self.assertEqual(arr.num_times, 0) self.assertEqual(arr.num_bytes, 4) self.assertEqual(arr.width, 0) @@ -130,7 +130,7 @@ def test_compute_scale_params_from_image_vect(self): def test_fill_search_data(self): for num_bytes in [2, 4]: - arr = SearchData() + arr = PsiPhiArray() fill_search_data( arr, num_bytes, [self.psi_1, self.psi_2], [self.phi_1, self.phi_2], self.zeroed_times, False ) @@ -191,8 +191,8 @@ def test_fill_search_data_from_image_stack(self): ) im_stack = ImageStack(images) - # Create the SearchData from the ImageStack. - arr = SearchData() + # Create the PsiPhiArray from the ImageStack. + arr = PsiPhiArray() fill_search_data_from_image_stack(arr, im_stack, 4, False) # Check the meta data. From a0605b6db638729a78904fb841b1cf52a1f5ca28 Mon Sep 17 00:00:00 2001 From: Jeremy Kubica <104161096+jeremykubica@users.noreply.github.com> Date: Fri, 2 Feb 2024 16:23:41 -0500 Subject: [PATCH 07/11] Update psi_phi_array.cpp --- src/kbmod/search/psi_phi_array.cpp | 112 ++++++++++++++--------------- 1 file changed, 56 insertions(+), 56 deletions(-) diff --git a/src/kbmod/search/psi_phi_array.cpp b/src/kbmod/search/psi_phi_array.cpp index ae2e1c42..329d7493 100644 --- a/src/kbmod/search/psi_phi_array.cpp +++ b/src/kbmod/search/psi_phi_array.cpp @@ -1,25 +1,25 @@ -#include "search_data_ds.h" -#include "search_data_utils.h" -#include "pydocs/search_data_docs.h" +#include "psi_phi_array_ds.h" +#include "psi_phi_array_utils.h" +#include "pydocs/psi_phi_array_docs.h" namespace search { // Declaration of CUDA functions that will be linked in. #ifdef HAVE_CUDA -extern "C" void device_allocate_search_data_arrays(SearchData* data); +extern "C" void device_allocate_psi_phi_array_arrays(PsiPhiArray* data); -extern "C" void device_free_search_data_arrays(SearchData* data); +extern "C" void device_free_psi_phi_array_arrays(PsiPhiArray* data); #endif // ------------------------------------------------------- // --- Implementation of core data structure functions --- // ------------------------------------------------------- -SearchData::SearchData() {} +PsiPhiArray::PsiPhiArray() {} -SearchData::~SearchData() { clear(); } +PsiPhiArray::~PsiPhiArray() { clear(); } -void SearchData::clear() { +void PsiPhiArray::clear() { // Free all used memory on CPU and GPU. if (cpu_array_ptr != nullptr) { free(cpu_array_ptr); @@ -31,7 +31,7 @@ void SearchData::clear() { } #ifdef HAVE_CUDA if ((gpu_array_ptr != nullptr) || (gpu_time_array != nullptr)) { - device_free_search_data_arrays(this); + device_free_psi_phi_array_arrays(this); gpu_array_ptr = nullptr; gpu_time_array = nullptr; } @@ -54,7 +54,7 @@ void SearchData::clear() { meta_data.phi_scale = 1.0; } -void SearchData::set_meta_data(int new_num_bytes, int new_num_times, int new_height, int new_width) { +void PsiPhiArray::set_meta_data(int new_num_bytes, int new_num_times, int new_height, int new_width) { // Validity checking of parameters. if (new_num_bytes != -1 && new_num_bytes != 1 && new_num_bytes != 2 && new_num_bytes != 4) { throw std::runtime_error("Invalid setting of num_bytes. Must be (-1 [use default], 1, 2, or 4)."); @@ -86,7 +86,7 @@ void SearchData::set_meta_data(int new_num_bytes, int new_num_times, int new_hei meta_data.total_array_size = meta_data.block_size * meta_data.num_entries; } -void SearchData::set_psi_scaling(float min_val, float max_val, float scale_val) { +void PsiPhiArray::set_psi_scaling(float min_val, float max_val, float scale_val) { if (min_val > max_val) throw std::runtime_error("Min value needs to be < max value"); if (scale_val <= 0) throw std::runtime_error("Scale value must be greater than zero."); meta_data.psi_min_val = min_val; @@ -94,7 +94,7 @@ void SearchData::set_psi_scaling(float min_val, float max_val, float scale_val) meta_data.psi_scale = scale_val; } -void SearchData::set_phi_scaling(float min_val, float max_val, float scale_val) { +void PsiPhiArray::set_phi_scaling(float min_val, float max_val, float scale_val) { if (min_val > max_val) throw std::runtime_error("Min value needs to be < max value"); if (scale_val <= 0) throw std::runtime_error("Scale value must be greater than zero."); meta_data.phi_min_val = min_val; @@ -102,7 +102,7 @@ void SearchData::set_phi_scaling(float min_val, float max_val, float scale_val) meta_data.phi_scale = scale_val; } -PsiPhi SearchData::read_psi_phi(int time, int row, int col) { +PsiPhi PsiPhiArray::read_psi_phi(int time, int row, int col) { PsiPhi result = {NO_DATA, NO_DATA}; // Array allocation and bounds checking. @@ -136,7 +136,7 @@ PsiPhi SearchData::read_psi_phi(int time, int row, int col) { return result; } -float SearchData::read_time(int time_index) { +float PsiPhiArray::read_time(int time_index) { if (cpu_time_array == nullptr) throw std::runtime_error("Read from unallocated times array."); if ((time_index < 0) || (time_index >= meta_data.num_times)) { throw std::runtime_error("Out of bounds read for time step."); @@ -175,8 +175,8 @@ std::array compute_scale_params_from_image_vect(const std::vector -void set_encode_cpu_search_data(SearchData& data, const std::vector& psi_imgs, - const std::vector& phi_imgs, bool debug) { +void set_encode_cpu_psi_phi_array(PsiPhiArray& data, const std::vector& psi_imgs, + const std::vector& phi_imgs, bool debug) { if (data.get_cpu_array_ptr() != nullptr) { throw std::runtime_error("CPU PsiPhi already allocated."); } @@ -219,8 +219,8 @@ void set_encode_cpu_search_data(SearchData& data, const std::vector& p data.set_cpu_array_ptr((void*)encoded); } -void set_float_cpu_search_data(SearchData& data, const std::vector& psi_imgs, - const std::vector& phi_imgs, bool debug) { +void set_float_cpu_psi_phi_array(PsiPhiArray& data, const std::vector& psi_imgs, + const std::vector& phi_imgs, bool debug) { if (data.get_cpu_array_ptr() != nullptr) { throw std::runtime_error("CPU PsiPhi already allocated."); } @@ -245,9 +245,9 @@ void set_float_cpu_search_data(SearchData& data, const std::vector& ps data.set_cpu_array_ptr((void*)encoded); } -void fill_search_data(SearchData& result_data, int num_bytes, const std::vector& psi_imgs, - const std::vector& phi_imgs, const std::vector zeroed_times, - bool debug) { +void fill_psi_phi_array(PsiPhiArray& result_data, int num_bytes, const std::vector& psi_imgs, + const std::vector& phi_imgs, const std::vector zeroed_times, + bool debug) { if (result_data.get_cpu_array_ptr() != nullptr) { return; } @@ -282,16 +282,16 @@ void fill_search_data(SearchData& result_data, int num_bytes, const std::vector< // Do the local encoding. if (result_data.get_num_bytes() == 1) { - set_encode_cpu_search_data(result_data, psi_imgs, phi_imgs, debug); + set_encode_cpu_psi_phi_array(result_data, psi_imgs, phi_imgs, debug); } else { - set_encode_cpu_search_data(result_data, psi_imgs, phi_imgs, debug); + set_encode_cpu_psi_phi_array(result_data, psi_imgs, phi_imgs, debug); } } else { if (debug) { printf("Encoding psi and phi as floats.\n"); } // Just interleave psi and phi images. - set_float_cpu_search_data(result_data, psi_imgs, phi_imgs, debug); + set_float_cpu_psi_phi_array(result_data, psi_imgs, phi_imgs, debug); } // Copy the time array. @@ -313,7 +313,7 @@ void fill_search_data(SearchData& result_data, int num_bytes, const std::vector< printf("Allocating GPU memory for times array using %lu bytes.\n", times_bytes); } - device_allocate_search_data_arrays(&result_data); + device_allocate_psi_phi_array_arrays(&result_data); if (result_data.get_gpu_array_ptr() == nullptr) { throw std::runtime_error("Unable to allocate GPU PsiPhi array."); } @@ -323,8 +323,8 @@ void fill_search_data(SearchData& result_data, int num_bytes, const std::vector< #endif } -void fill_search_data_from_image_stack(SearchData& result_data, ImageStack& stack, int num_bytes, - bool debug) { +void fill_psi_phi_array_from_image_stack(PsiPhiArray& result_data, ImageStack& stack, int num_bytes, + bool debug) { // Compute Phi and Psi from convolved images while leaving masked pixels alone // Reinsert 0s for NO_DATA? std::vector psi_images; @@ -346,7 +346,7 @@ void fill_search_data_from_image_stack(SearchData& result_data, ImageStack& stac // Convert these into an array form. Needs the full psi and phi computed first so the // encoding can compute the bounds of each array. std::vector zeroed_times = stack.build_zeroed_times(); - fill_search_data(result_data, num_bytes, psi_images, phi_images, zeroed_times, debug); + fill_psi_phi_array(result_data, num_bytes, psi_images, phi_images, zeroed_times, debug); } // ------------------------------------------- @@ -354,55 +354,55 @@ void fill_search_data_from_image_stack(SearchData& result_data, ImageStack& stac // ------------------------------------------- #ifdef Py_PYTHON_H -static void search_data_binding(py::module& m) { - using ppa = search::SearchData; +static void psi_phi_array_binding(py::module& m) { + using ppa = search::PsiPhiArray; py::class_(m, "PsiPhi", pydocs::DOC_PsiPhi) .def(py::init<>()) .def_readwrite("psi", &search::PsiPhi::psi) .def_readwrite("phi", &search::PsiPhi::phi); - py::class_(m, "SearchData", pydocs::DOC_SearchData) + py::class_(m, "PsiPhiArray", pydocs::DOC_PsiPhiArray) .def(py::init<>()) - .def_property_readonly("num_bytes", &ppa::get_num_bytes, pydocs::DOC_SearchData_get_num_bytes) - .def_property_readonly("num_times", &ppa::get_num_times, pydocs::DOC_SearchData_get_num_times) - .def_property_readonly("width", &ppa::get_width, pydocs::DOC_SearchData_get_width) - .def_property_readonly("height", &ppa::get_height, pydocs::DOC_SearchData_get_height) + .def_property_readonly("num_bytes", &ppa::get_num_bytes, pydocs::DOC_PsiPhiArray_get_num_bytes) + .def_property_readonly("num_times", &ppa::get_num_times, pydocs::DOC_PsiPhiArray_get_num_times) + .def_property_readonly("width", &ppa::get_width, pydocs::DOC_PsiPhiArray_get_width) + .def_property_readonly("height", &ppa::get_height, pydocs::DOC_PsiPhiArray_get_height) .def_property_readonly("pixels_per_image", &ppa::get_pixels_per_image, - pydocs::DOC_SearchData_get_pixels_per_image) + pydocs::DOC_PsiPhiArray_get_pixels_per_image) .def_property_readonly("num_entries", &ppa::get_num_entries, - pydocs::DOC_SearchData_get_num_entries) + pydocs::DOC_PsiPhiArray_get_num_entries) .def_property_readonly("total_array_size", &ppa::get_total_array_size, - pydocs::DOC_SearchData_get_total_array_size) - .def_property_readonly("block_size", &ppa::get_block_size, pydocs::DOC_SearchData_get_block_size) + pydocs::DOC_PsiPhiArray_get_total_array_size) + .def_property_readonly("block_size", &ppa::get_block_size, pydocs::DOC_PsiPhiArray_get_block_size) .def_property_readonly("psi_min_val", &ppa::get_psi_min_val, - pydocs::DOC_SearchData_get_psi_min_val) + pydocs::DOC_PsiPhiArray_get_psi_min_val) .def_property_readonly("psi_max_val", &ppa::get_psi_max_val, - pydocs::DOC_SearchData_get_psi_max_val) - .def_property_readonly("psi_scale", &ppa::get_psi_scale, pydocs::DOC_SearchData_get_psi_scale) + pydocs::DOC_PsiPhiArray_get_psi_max_val) + .def_property_readonly("psi_scale", &ppa::get_psi_scale, pydocs::DOC_PsiPhiArray_get_psi_scale) .def_property_readonly("phi_min_val", &ppa::get_phi_min_val, - pydocs::DOC_SearchData_get_phi_min_val) + pydocs::DOC_PsiPhiArray_get_phi_min_val) .def_property_readonly("phi_max_val", &ppa::get_phi_max_val, - pydocs::DOC_SearchData_get_phi_max_val) - .def_property_readonly("phi_scale", &ppa::get_phi_scale, pydocs::DOC_SearchData_get_phi_scale) + pydocs::DOC_PsiPhiArray_get_phi_max_val) + .def_property_readonly("phi_scale", &ppa::get_phi_scale, pydocs::DOC_PsiPhiArray_get_phi_scale) .def_property_readonly("cpu_array_allocated", &ppa::cpu_array_allocated, - pydocs::DOC_SearchData_get_cpu_array_allocated) + pydocs::DOC_PsiPhiArray_get_cpu_array_allocated) .def_property_readonly("gpu_array_allocated", &ppa::gpu_array_allocated, - pydocs::DOC_SearchData_get_gpu_array_allocated) + pydocs::DOC_PsiPhiArray_get_gpu_array_allocated) .def_property_readonly("cpu_time_array_allocated", &ppa::cpu_time_array_allocated, - pydocs::DOC_SearchData_get_cpu_time_array_allocated) + pydocs::DOC_PsiPhiArray_get_cpu_time_array_allocated) .def_property_readonly("gpu_time_array_allocated", &ppa::gpu_time_array_allocated, - pydocs::DOC_SearchData_get_gpu_time_array_allocated) - .def("set_meta_data", &ppa::set_meta_data, pydocs::DOC_SearchData_set_meta_data) - .def("clear", &ppa::clear, pydocs::DOC_SearchData_clear) - .def("read_psi_phi", &ppa::read_psi_phi, pydocs::DOC_SearchData_read_psi_phi) - .def("read_time", &ppa::read_time, pydocs::DOC_SearchData_read_time); + pydocs::DOC_PsiPhiArray_get_gpu_time_array_allocated) + .def("set_meta_data", &ppa::set_meta_data, pydocs::DOC_PsiPhiArray_set_meta_data) + .def("clear", &ppa::clear, pydocs::DOC_PsiPhiArray_clear) + .def("read_psi_phi", &ppa::read_psi_phi, pydocs::DOC_PsiPhiArray_read_psi_phi) + .def("read_time", &ppa::read_time, pydocs::DOC_PsiPhiArray_read_time); m.def("compute_scale_params_from_image_vect", &search::compute_scale_params_from_image_vect); m.def("decode_uint_scalar", &search::decode_uint_scalar); m.def("encode_uint_scalar", &search::encode_uint_scalar); - m.def("fill_search_data", &search::fill_search_data, pydocs::DOC_SearchData_fill_search_data); - m.def("fill_search_data_from_image_stack", &search::fill_search_data_from_image_stack, - pydocs::DOC_SearchData_fill_search_data_from_image_stack); + m.def("fill_psi_phi_array", &search::fill_psi_phi_array, pydocs::DOC_PsiPhiArray_fill_psi_phi_array); + m.def("fill_psi_phi_array_from_image_stack", &search::fill_psi_phi_array_from_image_stack, + pydocs::DOC_PsiPhiArray_fill_psi_phi_array_from_image_stack); } #endif From 089778242689439f88c59155c98082f91837c192 Mon Sep 17 00:00:00 2001 From: Jeremy Kubica <104161096+jeremykubica@users.noreply.github.com> Date: Fri, 2 Feb 2024 16:24:43 -0500 Subject: [PATCH 08/11] Rename file --- .../search/pydocs/{search_data_docs.h => psi_phi_array_docs.h} | 0 1 file changed, 0 insertions(+), 0 deletions(-) rename src/kbmod/search/pydocs/{search_data_docs.h => psi_phi_array_docs.h} (100%) diff --git a/src/kbmod/search/pydocs/search_data_docs.h b/src/kbmod/search/pydocs/psi_phi_array_docs.h similarity index 100% rename from src/kbmod/search/pydocs/search_data_docs.h rename to src/kbmod/search/pydocs/psi_phi_array_docs.h From f0062d9adc89b85e20c6204cf16d0527ebc7e402 Mon Sep 17 00:00:00 2001 From: Jeremy Kubica <104161096+jeremykubica@users.noreply.github.com> Date: Fri, 2 Feb 2024 16:27:14 -0500 Subject: [PATCH 09/11] Fix names from bad replace --- src/kbmod/search/kernels.cu | 4 ++-- src/kbmod/search/psi_phi_array.cpp | 4 ++-- 2 files changed, 4 insertions(+), 4 deletions(-) diff --git a/src/kbmod/search/kernels.cu b/src/kbmod/search/kernels.cu index c400c438..50c49775 100644 --- a/src/kbmod/search/kernels.cu +++ b/src/kbmod/search/kernels.cu @@ -23,7 +23,7 @@ namespace search { -extern "C" void device_allocate_psi_phi_array_arrays(PsiPhiArray *data) { +extern "C" void device_allocate_psi_phi_arrays(PsiPhiArray *data) { if (!data->cpu_array_allocated() || !data->cpu_time_array_allocated()) { throw std::runtime_error("CPU data is not allocated."); } @@ -47,7 +47,7 @@ extern "C" void device_allocate_psi_phi_array_arrays(PsiPhiArray *data) { data->set_gpu_time_array_ptr(device_times_ptr); } -extern "C" void device_free_psi_phi_array_arrays(PsiPhiArray *data) { +extern "C" void device_free_psi_phi_arrays(PsiPhiArray *data) { if (data->gpu_array_allocated()) { checkCudaErrors(cudaFree(data->get_gpu_array_ptr())); data->set_gpu_array_ptr(nullptr); diff --git a/src/kbmod/search/psi_phi_array.cpp b/src/kbmod/search/psi_phi_array.cpp index 329d7493..a7bad84e 100644 --- a/src/kbmod/search/psi_phi_array.cpp +++ b/src/kbmod/search/psi_phi_array.cpp @@ -6,9 +6,9 @@ namespace search { // Declaration of CUDA functions that will be linked in. #ifdef HAVE_CUDA -extern "C" void device_allocate_psi_phi_array_arrays(PsiPhiArray* data); +extern "C" void device_allocate_psi_phi_arrays(PsiPhiArray* data); -extern "C" void device_free_psi_phi_array_arrays(PsiPhiArray* data); +extern "C" void device_free_psi_phi_arrays(PsiPhiArray* data); #endif // ------------------------------------------------------- From 5d1d76cd4dce4cffec362f040b6e367987d33cb0 Mon Sep 17 00:00:00 2001 From: Jeremy Kubica <104161096+jeremykubica@users.noreply.github.com> Date: Fri, 2 Feb 2024 16:28:37 -0500 Subject: [PATCH 10/11] Update test_psi_phi_array.py --- tests/test_psi_phi_array.py | 14 +++++++------- 1 file changed, 7 insertions(+), 7 deletions(-) diff --git a/tests/test_psi_phi_array.py b/tests/test_psi_phi_array.py index 0e805dd7..c261e4ed 100644 --- a/tests/test_psi_phi_array.py +++ b/tests/test_psi_phi_array.py @@ -14,12 +14,12 @@ compute_scale_params_from_image_vect, decode_uint_scalar, encode_uint_scalar, - fill_search_data, - fill_search_data_from_image_stack, + fill_psi_phi_array, + fill_psi_phi_array_from_image_stack, ) -class test_search_data(unittest.TestCase): +class test_psi_phi_array(unittest.TestCase): def setUp(self): self.num_times = 2 self.width = 4 @@ -128,10 +128,10 @@ def test_compute_scale_params_from_image_vect(self): self.assertAlmostEqual(result_uint16[1], max_val, delta=1e-5) self.assertAlmostEqual(result_uint16[2], max_val / 65535.0, delta=1e-5) - def test_fill_search_data(self): + def test_fill_psi_phi_array(self): for num_bytes in [2, 4]: arr = PsiPhiArray() - fill_search_data( + fill_psi_phi_array( arr, num_bytes, [self.psi_1, self.psi_2], [self.phi_1, self.phi_2], self.zeroed_times, False ) @@ -173,7 +173,7 @@ def test_fill_search_data(self): self.assertFalse(arr.gpu_array_allocated) self.assertFalse(arr.gpu_time_array_allocated) - def test_fill_search_data_from_image_stack(self): + def test_fill_psi_phi_array_from_image_stack(self): # Build a fake image stack. num_times = 5 width = 21 @@ -193,7 +193,7 @@ def test_fill_search_data_from_image_stack(self): # Create the PsiPhiArray from the ImageStack. arr = PsiPhiArray() - fill_search_data_from_image_stack(arr, im_stack, 4, False) + fill_psi_phi_array_from_image_stack(arr, im_stack, 4, False) # Check the meta data. self.assertEqual(arr.num_times, num_times) From 17105527ffb9462af17c2b4c76b9e6d48f884b2b Mon Sep 17 00:00:00 2001 From: Jeremy Kubica <104161096+jeremykubica@users.noreply.github.com> Date: Fri, 2 Feb 2024 16:31:04 -0500 Subject: [PATCH 11/11] Update psi_phi_array.cpp --- src/kbmod/search/psi_phi_array.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/src/kbmod/search/psi_phi_array.cpp b/src/kbmod/search/psi_phi_array.cpp index a7bad84e..e0640e9d 100644 --- a/src/kbmod/search/psi_phi_array.cpp +++ b/src/kbmod/search/psi_phi_array.cpp @@ -31,7 +31,7 @@ void PsiPhiArray::clear() { } #ifdef HAVE_CUDA if ((gpu_array_ptr != nullptr) || (gpu_time_array != nullptr)) { - device_free_psi_phi_array_arrays(this); + device_free_psi_phi_arrays(this); gpu_array_ptr = nullptr; gpu_time_array = nullptr; } @@ -313,7 +313,7 @@ void fill_psi_phi_array(PsiPhiArray& result_data, int num_bytes, const std::vect printf("Allocating GPU memory for times array using %lu bytes.\n", times_bytes); } - device_allocate_psi_phi_array_arrays(&result_data); + device_allocate_psi_phi_arrays(&result_data); if (result_data.get_gpu_array_ptr() == nullptr) { throw std::runtime_error("Unable to allocate GPU PsiPhi array."); }