Improve pulscan integration and GPU robustness - #284
Closed
dejay-vu wants to merge 7 commits into
Closed
Conversation
- Introduced new source file `aa_device_pulscan.cu` containing CUDA kernels for Pulscan processing. - Updated `aa_pipeline_runner.hpp` to include Pulscan candidate retrieval method, conditionally compiled based on `AA_WITH_PULSCAN`. - Modified `meson.build` to add compilation flag for Pulscan support. - Included `aa_device_pulscan.cu` in the build source list. - Implemented various kernels for data processing, normalization, and candidate searching specific to Pulscan.
- Implemented test for calculating logp in `test_pulscan_calculate_logp.cpp` and corresponding CMake configuration. - Added sorting test for pulscan candidates in `test_pulscan_candidate_sort.cpp` with CMake setup. - Created decimation harmonics test in `test_pulscan_decimate_harmonics.cpp` and its CMake configuration. - Developed helper functions test in `test_pulscan_helpers.cpp` with appropriate CMake integration. - Introduced magnitude squared test in `test_pulscan_magnitude_squared.cpp` and its CMake setup. - Added median normalization test in `test_pulscan_median_normalisation.cpp` with CMake configuration. - Integrated pipeline testing in `test_pulscan_pipeline_integration.cpp` and corresponding CMake setup. - Created separate components test in `test_pulscan_separate_components.cpp` with CMake configuration.
## Summary This PR finishes the current Pulscan integration work in AstroAccelerate and tightens a few GPU-side failure modes that were uncovered during delivery testing. ## What changed - standardize the Pulscan build flag on `AA_ENABLE_PULSCAN` across CMake, Meson, Makefile, and docs - add the dedicated Pulscan runner and wire Pulscan output generation into the main pipeline flow - fix the Pulscan boxcar tail-handling bug that could emit bogus large candidates on short inputs - propagate `failsafe` correctly and harden DDTR/corner-turn memory and launch handling - add regression coverage for boxcar filtering, compatibility helpers, and the requirement that Pulscan only runs with periodicity enabled - update the README with build, install, runtime, output, and `AA_PULSCAN_DUMP_FFT` documentation - add helper configs and the filterbank split script used for validation ## Validation - rebuilt `astro-accelerate` successfully - ran the short repro `.fil` and confirmed the abnormal large Pulscan values disappeared - ran `build-tests/tests/test_pulscan_boxcar_filter` - ran the larger `73df` delivery input through AA and produced Pulscan output successfully
This reverts commit 576baaf.
There was a problem hiding this comment.
Pull request overview
This PR completes the in-tree Pulscan integration (build-flag standardization, runner wiring, and output emission) and adds GPU robustness improvements (notably corner-turn grid-y chunking), along with a suite of Pulscan-focused regression tests and helper scripts/configs.
Changes:
- Standardize and propagate
AA_ENABLE_PULSCANacross CMake/Meson/Make and conditionalize Pulscan tests and APIs. - Add Pulscan GPU kernels +
aa_pulscan_runnerand wire Pulscan outputs into the main executable and pipeline interfaces. - Improve GPU robustness for large launches (corner turn) and add validation helpers/tests plus sample configs/scripts.
Reviewed changes
Copilot reviewed 52 out of 53 changed files in this pull request and generated 7 comments.
Show a summary per file
| File | Description |
|---|---|
| tests/test_pulscan_separate_components.cpp | Adds CUDA unit test for complex-to-real/imag split kernel. |
| tests/test_pulscan_separate_components.cmake | Registers the separate-components test in CTest. |
| tests/test_pulscan_requires_periodicity.cpp | Verifies Pulscan does not produce candidates without periodicity enabled. |
| tests/test_pulscan_requires_periodicity.cmake | Registers requires-periodicity test in CTest. |
| tests/test_pulscan_pipeline_integration.cpp | Runs a minimal pipeline with periodicity and checks candidate ordering. |
| tests/test_pulscan_pipeline_integration.cmake | Registers pipeline integration test in CTest. |
| tests/test_pulscan_median_normalisation.cpp | Adds CUDA unit test for median/MAD normalization kernel. |
| tests/test_pulscan_median_normalisation.cmake | Registers median-normalisation test in CTest. |
| tests/test_pulscan_magnitude_squared.cpp | Adds CUDA unit test for magnitude-squared kernel. |
| tests/test_pulscan_magnitude_squared.cmake | Registers magnitude-squared test in CTest. |
| tests/test_pulscan_helpers.cpp | Tests host helper pulscan_logp_to_sigma. |
| tests/test_pulscan_helpers.cmake | Registers helpers test in CTest. |
| tests/test_pulscan_decimate_harmonics.cpp | Adds CUDA unit test for harmonic decimation kernel. |
| tests/test_pulscan_decimate_harmonics.cmake | Registers decimate-harmonics test in CTest. |
| tests/test_pulscan_compatibility_helpers.cpp | Tests FFT-bin truncation helper. |
| tests/test_pulscan_compatibility_helpers.cmake | Registers compatibility helpers test in CTest. |
| tests/test_pulscan_candidate_sort.cpp | Tests candidate sort/trim logic on host. |
| tests/test_pulscan_candidate_sort.cmake | Registers candidate-sort test in CTest. |
| tests/test_pulscan_calculate_logp.cpp | Adds CUDA unit test for logp calculation kernel. |
| tests/test_pulscan_calculate_logp.cmake | Registers calculate-logp test in CTest. |
| tests/test_pulscan_boxcar_filter.cpp | Adds regression test for boxcar tail-handling + candidate emission. |
| tests/test_pulscan_boxcar_filter.cmake | Registers boxcar-filter test in CTest. |
| tests/CMakeLists.txt | Conditionally includes Pulscan tests behind AA_ENABLE_PULSCAN. |
| src/aa_main.cpp | Adds Pulscan output directory handling and writes Pulscan outputs after periodicity. |
| src/aa_device_pulscan.cu | Introduces Pulscan CUDA kernels + kernel-call wrappers + sigma helper. |
| src/aa_device_corner_turn_kernel.cu | Adds block_y_offset to support chunked launches in Y dimension. |
| src/aa_corner_turn.cpp | Chunks corner-turn launches to avoid exceeding CUDA grid-y limits. |
| scripts/split_filterbank.py | Adds helper script to split SIGPROC filterbanks into chunks. |
| meson_srclist.txt | Adds Pulscan CUDA source to Meson build source list. |
| meson.build | Adds Meson option mapping to AA_ENABLE_PULSCAN. |
| input_files/ska_test_file.txt | Updates sample config (currently with a local absolute path). |
| input_files/pulscan_test.cfg | Adds Pulscan validation config (contains absolute path). |
| input_files/pulscan_sample.cfg | Adds Pulscan sample config (contains absolute path). |
| input_files/pulscan_repro_no_periodicity.cfg | Adds repro config for “no periodicity” case (contains absolute path). |
| input_files/pulscan_gpu.cu | Adds standalone Pulscan validation CUDA program reference source. |
| input_files/pulscan_generated.cfg | Adds generated-data Pulscan config (contains absolute path). |
| input_files/pulscan_dm30125_part000_delivery.cfg | Adds delivery-validation config (contains absolute paths). |
| input_files/pulscan_dm30125_part000.cfg | Adds delivery-validation config (contains absolute path). |
| input_files/pulscan_dm30125.cfg | Adds delivery-validation config (contains absolute path). |
| include/aa_pulscan_runner.hpp | Adds the Pulscan runner (workspace mgmt, FFT, candidate generation, dumping). |
| include/aa_pipeline_runner.hpp | Extends pipeline runner interface with Pulscan getters/writers (guarded). |
| include/aa_pipeline_api.hpp | Exposes Pulscan getters/writers on the public API (guarded). |
| include/aa_pipeline.hpp | Adds failsafe pipeline option enum + description. |
| include/aa_params.hpp | Defines default AA_ENABLE_PULSCAN when not provided by build system. |
| include/aa_device_pulscan.hpp | Declares Pulscan structs and kernel wrapper APIs. |
| include/aa_device_info.hpp | Hardens device clock-rate handling using attributes for newer CUDART. |
| include/aa_device_corner_turn_kernel.hpp | Updates corner-turn kernel wrapper signatures to include block_y_offset. |
| include/aa_config.hpp | Adds Pulscan output directory flag parsing and wires failsafe option propagation. |
| README.md | Documents Pulscan build/runtime behavior and debug FFT dumping. |
| Makefile | Adds PULSCAN=1 mapping to -DAA_ENABLE_PULSCAN. |
| MANUAL.md | Adds Pulscan build/use notes (experimental). |
| CMakeLists.txt | Adds AA_ENABLE_PULSCAN option, updates C++ standard flags, adjusts linking visibility. |
| .gitignore | Ignores Pulscan output products. |
| build-tests/examples/CTestTestfile.cmake | Generated build artifact committed under build-tests/. |
| build-tests/astroaccelerate.pc | Generated build artifact committed under build-tests/. |
| build-tests/astro-accelerateConfig.cmake | Generated build artifact committed under build-tests/. |
| build-tests/Testing/Temporary/LastTestsFailed.log | Generated build artifact committed under build-tests/. |
| build-tests/Testing/Temporary/LastTest.log | Generated build artifact committed under build-tests/. |
| build-tests/DartConfiguration.tcl | Generated build artifact committed under build-tests/. |
| build-tests/CTestTestfile.cmake | Generated build artifact committed under build-tests/. |
💡 Add Copilot custom instructions for smarter, more guided reviews. Learn how to get started.
| copy_ddtr_data_to_host | ||
| threshold | ||
| failsafe | ||
| file /home/jay/workspace/aa/pulscan/sample_data/1298us_binary_multi32_8bit.fil |
There was a problem hiding this comment.
This sample config also hardcodes a developer-specific absolute .fil path. For repository examples/validation configs, use relative paths or placeholders so others can run them without editing the file.
Suggested change
| file /home/jay/workspace/aa/pulscan/sample_data/1298us_binary_multi32_8bit.fil | |
| file sample_data/1298us_binary_multi32_8bit.fil |
Comment on lines
+319
to
+333
| if (strcmp(string, "pulscan-output-dir") == 0 || | ||
| strcmp(string, "pulscan_output_dir") == 0) { | ||
| wordexp_t expanded_string; | ||
|
|
||
| if (fscanf(fp_in, "%s", string) == 0) | ||
| { | ||
| fprintf(stderr, "failed to read input\n"); | ||
| return false; | ||
| } | ||
| wordexp(string, &expanded_string, 0); | ||
| if(expanded_string.we_wordc > 0) { | ||
| flg.pulscan_output_dir = expanded_string.we_wordv[0]; | ||
| } | ||
| wordfree(&expanded_string); | ||
| } |
There was a problem hiding this comment.
wordexp() is called with flags 0 and the return status is not checked. This allows command substitution (e.g. $(...)) from the config file and can fail silently, which is a security risk for a user-supplied config. Use WRDE_NOCMD (and check the wordexp() return code) before trusting we_wordv, and handle errors by rejecting the config entry.
| sigma_cutoff 6 | ||
| periodicity | ||
| periodicity_sigma_cutoff 3 | ||
| periodicity_nHarmonics 32 |
There was a problem hiding this comment.
The README example uses periodicity_nHarmonics, but the config parser key is periodicity_harmonics (see aa_config.hpp). As written, the example won’t set the harmonics value. Update the snippet to use the supported key name.
Suggested change
| periodicity_nHarmonics 32 | |
| periodicity_harmonics 32 |
Comment on lines
+631
to
+636
| if(series_ptr == nullptr) { | ||
| continue; | ||
| } | ||
| std::copy_n(series_ptr, | ||
| host_stride, | ||
| host_buffer + static_cast<size_t>(dm_idx) * host_stride); |
There was a problem hiding this comment.
When series_ptr is null, the loop continues without writing anything into the corresponding portion of host_buffer. Since host_buffer is pinned memory from cudaMallocHost, it is not guaranteed to be zero-initialized, so uninitialized/stale data can be copied to the GPU and produce nondeterministic Pulscan output. Either explicitly zero-fill that DM slice, or skip launching Pulscan for that DM index (and adjust dm_series/batching accordingly).
Suggested change
| if(series_ptr == nullptr) { | |
| continue; | |
| } | |
| std::copy_n(series_ptr, | |
| host_stride, | |
| host_buffer + static_cast<size_t>(dm_idx) * host_stride); | |
| float* host_series = | |
| host_buffer + static_cast<size_t>(dm_idx) * host_stride; | |
| if(series_ptr == nullptr) { | |
| std::fill_n(host_series, host_stride, 0.0f); | |
| continue; | |
| } | |
| std::copy_n(series_ptr, host_stride, host_series); |
Comment on lines
+754
to
+783
| cudaMemcpy(m_workspace.h_fft_real, | ||
| m_workspace.d_real, | ||
| real_bytes, | ||
| cudaMemcpyDeviceToHost); | ||
| cudaMemcpy(m_workspace.h_fft_imag, | ||
| m_workspace.d_imag, | ||
| real_bytes, | ||
| cudaMemcpyDeviceToHost); | ||
| const std::string dump_dir = | ||
| pulscan_fft_dump_directory(dump_setting); | ||
| if(!pulscan_ensure_directory_exists(dump_dir)) { | ||
| LOG(log_level::error, | ||
| "Pulscan: failed to create FFT dump directory " + dump_dir); | ||
| } | ||
| else { | ||
| char filename[256]; | ||
| std::snprintf(filename, | ||
| sizeof(filename), | ||
| "pulscan_fft_dm%03d_part%d.bin", | ||
| plan_entry.range_index, | ||
| dm_idx); | ||
| const std::string filepath = | ||
| pulscan_join_path(dump_dir, filename); | ||
| FILE* fp_fft = std::fopen(filepath.c_str(), "wb"); | ||
| if(fp_fft) { | ||
| for(size_t i = 0; i < usable_freq_bins; ++i) { | ||
| std::fwrite(&m_workspace.h_fft_real[i], sizeof(float), 1, fp_fft); | ||
| std::fwrite(&m_workspace.h_fft_imag[i], sizeof(float), 1, fp_fft); | ||
| } | ||
| std::fclose(fp_fft); |
There was a problem hiding this comment.
The FFT dump path ignores the return values of cudaMemcpy. If the device->host copies fail, the code will still write whatever is in the host buffers to disk, making dumps misleading and potentially hiding GPU errors. Capture and check the cudaMemcpy statuses (and consider early-out / logging on failure).
Suggested change
| cudaMemcpy(m_workspace.h_fft_real, | |
| m_workspace.d_real, | |
| real_bytes, | |
| cudaMemcpyDeviceToHost); | |
| cudaMemcpy(m_workspace.h_fft_imag, | |
| m_workspace.d_imag, | |
| real_bytes, | |
| cudaMemcpyDeviceToHost); | |
| const std::string dump_dir = | |
| pulscan_fft_dump_directory(dump_setting); | |
| if(!pulscan_ensure_directory_exists(dump_dir)) { | |
| LOG(log_level::error, | |
| "Pulscan: failed to create FFT dump directory " + dump_dir); | |
| } | |
| else { | |
| char filename[256]; | |
| std::snprintf(filename, | |
| sizeof(filename), | |
| "pulscan_fft_dm%03d_part%d.bin", | |
| plan_entry.range_index, | |
| dm_idx); | |
| const std::string filepath = | |
| pulscan_join_path(dump_dir, filename); | |
| FILE* fp_fft = std::fopen(filepath.c_str(), "wb"); | |
| if(fp_fft) { | |
| for(size_t i = 0; i < usable_freq_bins; ++i) { | |
| std::fwrite(&m_workspace.h_fft_real[i], sizeof(float), 1, fp_fft); | |
| std::fwrite(&m_workspace.h_fft_imag[i], sizeof(float), 1, fp_fft); | |
| } | |
| std::fclose(fp_fft); | |
| const cudaError_t real_copy_status = | |
| cudaMemcpy(m_workspace.h_fft_real, | |
| m_workspace.d_real, | |
| real_bytes, | |
| cudaMemcpyDeviceToHost); | |
| const cudaError_t imag_copy_status = | |
| cudaMemcpy(m_workspace.h_fft_imag, | |
| m_workspace.d_imag, | |
| real_bytes, | |
| cudaMemcpyDeviceToHost); | |
| if(real_copy_status != cudaSuccess || imag_copy_status != cudaSuccess) { | |
| std::string error_message = | |
| "Pulscan: failed to copy FFT dump buffers to host"; | |
| if(real_copy_status != cudaSuccess) { | |
| error_message += | |
| std::string(" (real: ") + cudaGetErrorString(real_copy_status); | |
| if(imag_copy_status != cudaSuccess) { | |
| error_message += | |
| std::string(", imag: ") + cudaGetErrorString(imag_copy_status); | |
| } | |
| error_message += ")"; | |
| } | |
| else { | |
| error_message += | |
| std::string(" (imag: ") + cudaGetErrorString(imag_copy_status) + | |
| ")"; | |
| } | |
| LOG(log_level::error, error_message); | |
| } | |
| else { | |
| const std::string dump_dir = | |
| pulscan_fft_dump_directory(dump_setting); | |
| if(!pulscan_ensure_directory_exists(dump_dir)) { | |
| LOG(log_level::error, | |
| "Pulscan: failed to create FFT dump directory " + dump_dir); | |
| } | |
| else { | |
| char filename[256]; | |
| std::snprintf(filename, | |
| sizeof(filename), | |
| "pulscan_fft_dm%03d_part%d.bin", | |
| plan_entry.range_index, | |
| dm_idx); | |
| const std::string filepath = | |
| pulscan_join_path(dump_dir, filename); | |
| FILE* fp_fft = std::fopen(filepath.c_str(), "wb"); | |
| if(fp_fft) { | |
| for(size_t i = 0; i < usable_freq_bins; ++i) { | |
| std::fwrite(&m_workspace.h_fft_real[i], sizeof(float), 1, fp_fft); | |
| std::fwrite(&m_workspace.h_fft_imag[i], sizeof(float), 1, fp_fft); | |
| } | |
| std::fclose(fp_fft); | |
| } |
Comment on lines
+40
to
+89
| bool ensure_directory_exists(const std::string& path) { | ||
| if(path.empty() || path == ".") { | ||
| return true; | ||
| } | ||
|
|
||
| std::string current; | ||
| size_t position = 0; | ||
| if(path[0] == '/') { | ||
| current = "/"; | ||
| position = 1; | ||
| } | ||
|
|
||
| while(position <= path.size()) { | ||
| const size_t next_separator = path.find('/', position); | ||
| const std::string part = path.substr(position, next_separator - position); | ||
| if(!part.empty() && part != ".") { | ||
| if(!current.empty() && current.back() != '/') { | ||
| current += "/"; | ||
| } | ||
| current += part; | ||
|
|
||
| struct stat path_stat; | ||
| if(stat(current.c_str(), &path_stat) != 0) { | ||
| if(mkdir(current.c_str(), 0755) != 0 && errno != EEXIST) { | ||
| return false; | ||
| } | ||
| } | ||
| else if(!S_ISDIR(path_stat.st_mode)) { | ||
| return false; | ||
| } | ||
| } | ||
|
|
||
| if(next_separator == std::string::npos) { | ||
| break; | ||
| } | ||
| position = next_separator + 1; | ||
| } | ||
|
|
||
| return true; | ||
| } | ||
|
|
||
| std::string join_path(const std::string& directory, const std::string& filename) { | ||
| if(directory.empty() || directory == ".") { | ||
| return filename; | ||
| } | ||
| if(directory.back() == '/') { | ||
| return directory + filename; | ||
| } | ||
| return directory + "/" + filename; | ||
| } |
There was a problem hiding this comment.
ensure_directory_exists() / join_path() duplicate the Pulscan path helpers that already exist in the Pulscan implementation (e.g. pulscan_ensure_directory_exists / pulscan_join_path). Duplicating this logic in aa_main.cpp risks divergence (different path handling / permissions / separators). Prefer reusing a single shared helper (e.g. move into a small aa_fs_utils header used by both sites).
| sigma_cutoff 6 | ||
| periodicity | ||
| dedispersion | ||
| file /home/jay/workspace/aa/pulscan/sample_data/1298us_binary_from_dat.fil |
There was a problem hiding this comment.
This config file hardcodes an absolute local path to a .fil input (/home/jay/...). That makes the sample unusable for other developers/CI. Prefer a relative path under the repo (or a placeholder like /path/to/file.fil) and add a brief comment explaining what the file should be.
Suggested change
| file /home/jay/workspace/aa/pulscan/sample_data/1298us_binary_from_dat.fil | |
| # Path to the sample .fil input file relative to the repository root. | |
| file sample_data/1298us_binary_from_dat.fil |
Collaborator
Author
|
Superseded by #286, which recreates the |
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Summary
Finish the current Pulscan integration work and ship the GPU robustness fixes into
master.Highlights
AA_ENABLE_PULSCANValidation
astro-accelerate.filand confirmed the abnormal large candidates disappearedbuild-tests/tests/test_pulscan_boxcar_filter