diff --git a/.github/workflows/e2e.yml b/.github/workflows/e2e.yml index 93799a9..6204469 100644 --- a/.github/workflows/e2e.yml +++ b/.github/workflows/e2e.yml @@ -61,6 +61,7 @@ jobs: ssh -o "StrictHostKeyChecking no" $INSTANCE_USERNAME@$instance_ip "cp -r /home/$INSTANCE_USERNAME/e2e /dev/shm/" - uses: actions/checkout@v4 - name: End-to-end tests + continue-on-error: true env: INSTANCE_USERNAME: ${{ secrets.EAGLE_USERNAME }} ACCESS_KEY_ID: ${{ secrets.PALMY_ACCESS_KEY_ID }} @@ -75,37 +76,19 @@ jobs: echo "aws_secret_access_key=${SECRET_ACCESS_KEY}" >> /tmp/cred scp -o "StrictHostKeyChecking no" /tmp/cred $INSTANCE_USERNAME@$instance_ip:/tmp/ ssh -o "StrictHostKeyChecking no" $INSTANCE_USERNAME@$instance_ip "/tmp/asar-focus/build-automation/run_e2e_container.sh /dev/shm/e2e $exe_image_name /tmp/cred" - # Separate step because when tests fail we would like to analyze the results still - stash-assets-from-ramdisk: - name: Stash E2E assets from ramdisk to persistent disk - runs-on: ubuntu-latest - needs: [start-gpu-instance, end-to-end-testing] - if: | - needs.end-to-end-testing.result != 'skipped' && - needs.end-to-end-testing.result != 'cancelled' - steps: - - name: Create SSH key - env: - SSH_PRIVATE_KEY: ${{ secrets.EAGLE_SSH_KEY }} - run: | - key_location=~/.ssh/ - mkdir -p $key_location - key_path=$key_location/id_rsa - echo "$SSH_PRIVATE_KEY" > $key_path - sudo chmod 600 $key_path - ssh-keygen -f $key_path -y > $key_path.pub - - name: Modify ownership and move assets + - name: Stash assets from ramdisk env: INSTANCE_USERNAME: ${{ secrets.EAGLE_USERNAME }} run: | instance_ip=${{ needs.start-gpu-instance.outputs.instance_ip }} ssh -o "StrictHostKeyChecking no" $INSTANCE_USERNAME@$instance_ip "sudo chown -R $INSTANCE_USERNAME:$INSTANCE_USERNAME /dev/shm/e2e" ssh -o "StrictHostKeyChecking no" $INSTANCE_USERNAME@$instance_ip "rsync -av /dev/shm/e2e /home/$INSTANCE_USERNAME/" + ssh -o "StrictHostKeyChecking no" $INSTANCE_USERNAME@$instance_ip "rm -rf /dev/shm/e2e" stop-gpu-instance: name: Stop GPU instance runs-on: ubuntu-latest - needs: [start-gpu-instance, end-to-end-testing, stash-assets-from-ramdisk] + needs: [start-gpu-instance, end-to-end-testing] if: | always() && needs.start-gpu-instance.result == 'success' diff --git a/CMakeLists.txt b/CMakeLists.txt index b7719f0..92a1340 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -23,18 +23,8 @@ set(SAR_SOURCES sar/sar_metadata.h ) -set(CUDA_UTIL_SOURCES - cuda_util/cuda_cleanup.h - cuda_util/cuda_util.h - cuda_util/cuda_workplace.h - cuda_util/cufft_plan.cpp - cuda_util/cufft_plan.h - cuda_util/device_padded_image.cu - cuda_util/device_padded_image.cuh - ) - set(SOURCES - main.cc main_flow.cc ${SAR_SOURCES} ${CUDA_UTIL_SOURCES} + main.cc main_flow.cc ${SAR_SOURCES} ) add_executable(asar_focus ${SOURCES}) @@ -76,13 +66,14 @@ include(FetchContent) include(dependencies/FetchContents.cmake) add_subdirectory(app) +add_subdirectory(cuda_util) add_subdirectory(envisat_format) add_subdirectory(util) target_include_directories(asar_focus PRIVATE ${CMAKE_CURRENT_LIST_DIR}) if (CMAKE_CXX_COMPILER_ID STREQUAL "GNU" AND CMAKE_CXX_COMPILER_VERSION VERSION_LESS 9) message("Also adding 'stdc++fs' for linking since the g++ version ${CMAKE_CXX_COMPILER_VERSION} requires it.") - target_link_libraries(asar_focus PRIVATE gdal cufft Eigen3::Eigen stdc++fs app-static util-static envisat-format-static) + target_link_libraries(asar_focus PRIVATE gdal cufft Eigen3::Eigen stdc++fs app-static util-static envisat-format-static cuda-util-static) else () - target_link_libraries(asar_focus PRIVATE gdal cufft Eigen3::Eigen app-static util-static envisat-format-static) + target_link_libraries(asar_focus PRIVATE gdal cufft Eigen3::Eigen app-static util-static envisat-format-static cuda-util-static) endif() diff --git a/build-automation/create_exe_in_container.sh b/build-automation/create_exe_in_container.sh index 6d8fc62..55e4bec 100755 --- a/build-automation/create_exe_in_container.sh +++ b/build-automation/create_exe_in_container.sh @@ -40,6 +40,8 @@ docker pull "${docker_image}" set +e docker stop "${container_name}" docker rm "${container_name}" +# Also remove images built before along with children. +docker rmi -f "${container_name}-exe" set -e docker run -t -d --name "${container_name}" "${docker_image}" docker cp "${repo_dir}" "${container_name}":"${container_work_dir}/" diff --git a/build-automation/run_e2e_container.sh b/build-automation/run_e2e_container.sh index cf43286..5cae805 100755 --- a/build-automation/run_e2e_container.sh +++ b/build-automation/run_e2e_container.sh @@ -20,7 +20,7 @@ if [ -d "$e2e_dir" ]; then # Extract the last folder name e2e_dir_name=$(basename "$e2e_dir") else - echo "Given repo dir ${e2e_dir} is not a directory or does not exist." + echo "Given E2E dir '${e2e_dir}' is not a directory or does not exist." exit 1 fi diff --git a/cuda_util/CMakeLists.txt b/cuda_util/CMakeLists.txt new file mode 100644 index 0000000..82f4490 --- /dev/null +++ b/cuda_util/CMakeLists.txt @@ -0,0 +1,25 @@ +list(APPEND CUDA_UTIL_INCLUDE_DIRS + ${CMAKE_CURRENT_LIST_DIR}) + +list(APPEND CUDA_UTIL_SOURCES + ${CMAKE_CURRENT_LIST_DIR}/cuda_device.cu + ${CMAKE_CURRENT_LIST_DIR}/cuda_device_init.cu + ${CMAKE_CURRENT_LIST_DIR}/cufft_plan.cc + ${CMAKE_CURRENT_LIST_DIR}/device_padded_image.cu + ${CMAKE_CURRENT_LIST_DIR}/memory_policy.cc +) + +add_library(cuda-util-static STATIC ${CUDA_UTIL_SOURCES}) +target_include_directories(cuda-util-static + PUBLIC + ${CUDA_UTIL_INCLUDE_DIRS} +) +target_link_libraries(cuda-util-static + PUBLIC + CUDA::cudart + app-static +) +set_target_properties(cuda-util-static + PROPERTIES + OUTPUT_NAME cuda-util +) diff --git a/cuda_util/cuda_device.cu b/cuda_util/cuda_device.cu new file mode 100644 index 0000000..fcb2976 --- /dev/null +++ b/cuda_util/cuda_device.cu @@ -0,0 +1,44 @@ +/** + * ENVISAT and ERS ASAR instrument focusser for QA4EO activity (c) by CGI Estonia AS + * + * ENVISAT and ERS ASAR instrument focusser for QA4EO activity is licensed under a + * Creative Commons Attribution-ShareAlike 4.0 International License. + * + * You should have received a copy of the license along with this + * work. If not, see http://creativecommons.org/licenses/by-sa/4.0/ + */ + +#include "cuda_device.h" + +#include + +#include "cuda_util.h" + +namespace alus::cuda { + +CudaDevice::CudaDevice(int device_nr, void* device_prop) : device_nr_{device_nr} { + cudaDeviceProp* dev = reinterpret_cast(device_prop); + cc_major_ = dev->major; + cc_minor_ = dev->minor; + name_ = dev->name; + sm_count_ = dev->multiProcessorCount; + max_threads_per_sm_ = dev->maxThreadsPerMultiProcessor; + warp_size_ = dev->warpSize; + total_global_memory_ = dev->totalGlobalMem; + alignment_ = dev->textureAlignment; +} + +void CudaDevice::Set() const { + CHECK_CUDA_ERR(cudaSetDevice(device_nr_)); +} + +size_t CudaDevice::GetFreeGlobalMemory() const { + Set(); + size_t total; + size_t free; + CHECK_CUDA_ERR(cudaMemGetInfo(&free, &total)); + + return free; +} + +} // namespace alus::cuda diff --git a/cuda_util/cuda_device.h b/cuda_util/cuda_device.h new file mode 100644 index 0000000..60b1735 --- /dev/null +++ b/cuda_util/cuda_device.h @@ -0,0 +1,51 @@ +/** +* ENVISAT and ERS ASAR instrument focusser for QA4EO activity (c) by CGI Estonia AS +* +* ENVISAT and ERS ASAR instrument focusser for QA4EO activity is licensed under a +* Creative Commons Attribution-ShareAlike 4.0 International License. +* +* You should have received a copy of the license along with this +* work. If not, see http://creativecommons.org/licenses/by-sa/4.0/ +*/ + +#pragma once + +#include +#include +#include + +namespace alus::cuda { +class CudaDevice final { +public: + CudaDevice() = delete; + /* + * Parses GPU device properties from 'cudaDeviceProp' struct pointer. + * It is opaque one here in order to not include CUDA SDK headers to host compilation. + */ + CudaDevice(int device_nr, void* device_prop); + + void Set() const; + + [[nodiscard]] int GetDeviceNr() const { return device_nr_; } + [[nodiscard]] std::string_view GetName() const { return name_; } + [[nodiscard]] size_t GetCcMajor() const { return cc_major_; } + [[nodiscard]] size_t GetCcMinor() const { return cc_minor_; } + [[nodiscard]] size_t GetSmCount() const { return sm_count_; } + [[nodiscard]] size_t GetMaxThreadsPerSm() const { return max_threads_per_sm_; } + [[nodiscard]] size_t GetWarpSize() const { return warp_size_; } + [[nodiscard]] size_t GetTotalGlobalMemory() const { return total_global_memory_; }; + [[nodiscard]] size_t GetFreeGlobalMemory() const; + [[nodiscard]] size_t GetMemoryAlignment() const { return alignment_; } + +private: + int device_nr_; + size_t cc_major_; + size_t cc_minor_; + std::string name_; + size_t sm_count_; + size_t max_threads_per_sm_; + size_t warp_size_; + size_t total_global_memory_; + size_t alignment_; +}; +} // namespace alus::cuda diff --git a/cuda_util/cuda_device_init.cu b/cuda_util/cuda_device_init.cu new file mode 100644 index 0000000..178cc77 --- /dev/null +++ b/cuda_util/cuda_device_init.cu @@ -0,0 +1,73 @@ +/** +* ENVISAT and ERS ASAR instrument focusser for QA4EO activity (c) by CGI Estonia AS +* +* ENVISAT and ERS ASAR instrument focusser for QA4EO activity is licensed under a +* Creative Commons Attribution-ShareAlike 4.0 International License. +* +* You should have received a copy of the license along with this +* work. If not, see http://creativecommons.org/licenses/by-sa/4.0/ +*/ + +#include "cuda_device_init.h" + +#include +#include + +#include + +#include "cuda_util.h" + +namespace alus::cuda { +CudaInit::CudaInit() { + init_future_ = std::async(std::launch::async, [this]() { this->QueryDevices(); }); +} + +bool CudaInit::IsFinished() const { + if (!init_future_.valid()) { + throw std::runtime_error("The future is already a past, invalid state queried."); + } + return init_future_.wait_for(std::chrono::milliseconds(0)) == std::future_status::ready; +} + +void CudaInit::QueryDevices() { + int device_count{}; + CHECK_CUDA_ERR(cudaGetDeviceCount(&device_count)); + if (!device_count) { + throw std::runtime_error("No GPU devices detected"); + } + for (int i{}; i < device_count; i++) { + cudaDeviceProp deviceProp; + CHECK_CUDA_ERR(cudaGetDeviceProperties(&deviceProp, i)); + devices_.emplace_back(i, &deviceProp); + // Whatever will first start invoking GPU, might be delayed if this thread does not finish. + // But when waiting, a first invocation of GPU could be delayed by waiting here. + // Also no error checking is done, because if there are errors, then sooner or later they will pop out + // somewhere else. + device_warmups_.emplace_back([i]() { + cudaSetDevice(i); + cudaFree(nullptr); + }); + } +} + +void CudaInit::CheckErrors() { + if (!init_future_.valid()) { + throw std::runtime_error("The future is already a past, invalid state queried."); + } + init_future_.get(); +} + +CudaInit::~CudaInit() { + // Just in case wait if any left hanging. + if (init_future_.valid()) { + init_future_.wait_for(std::chrono::seconds(10)); + } + + for (auto& t : device_warmups_) { + if (t.joinable()) { + t.join(); + } + } +} + +} // namespace alus::cuda \ No newline at end of file diff --git a/cuda_util/cuda_device_init.h b/cuda_util/cuda_device_init.h new file mode 100644 index 0000000..3162d1e --- /dev/null +++ b/cuda_util/cuda_device_init.h @@ -0,0 +1,38 @@ +/** +* ENVISAT and ERS ASAR instrument focusser for QA4EO activity (c) by CGI Estonia AS +* +* ENVISAT and ERS ASAR instrument focusser for QA4EO activity is licensed under a +* Creative Commons Attribution-ShareAlike 4.0 International License. +* +* You should have received a copy of the license along with this +* work. If not, see http://creativecommons.org/licenses/by-sa/4.0/ +*/ + +#pragma once + +#include +#include +#include + +#include "cuda_device.h" + +namespace alus::cuda { +class CudaInit final { +public: + CudaInit(); + + [[nodiscard]] bool IsFinished() const; + void CheckErrors(); + + [[nodiscard]] const std::vector& GetDevices() const { return devices_; } + + ~CudaInit(); + +private: + void QueryDevices(); + + std::vector devices_; + std::future init_future_; + std::vector device_warmups_; +}; +} // namespace alus::cuda \ No newline at end of file diff --git a/cuda_util/cuda_util.h b/cuda_util/cuda_util.h index 76cdd54..e89e3fd 100644 --- a/cuda_util/cuda_util.h +++ b/cuda_util/cuda_util.h @@ -7,6 +7,7 @@ * You should have received a copy of the license along with this * work. If not, see http://creativecommons.org/licenses/by-sa/4.0/ */ + #pragma once #include diff --git a/cuda_util/cufft_checks.h b/cuda_util/cufft_checks.h new file mode 100644 index 0000000..dd0a580 --- /dev/null +++ b/cuda_util/cufft_checks.h @@ -0,0 +1,93 @@ +/** +* ENVISAT and ERS ASAR instrument focusser for QA4EO activity (c) by CGI Estonia AS +* +* ENVISAT and ERS ASAR instrument focusser for QA4EO activity is licensed under a +* Creative Commons Attribution-ShareAlike 4.0 International License. +* +* You should have received a copy of the license along with this +* work. If not, see http://creativecommons.org/licenses/by-sa/4.0/ + */ + +#include + +#include + +#define CHECK_CUFFT_ERR(x) alus::cuda::CheckCuFFT(x, __FILE__, __LINE__) + +namespace alus::cuda { + +// copy paste from cuda-11.2/samples/common/inc/helper_cuda.h +inline const char* CufftErrorStr(cufftResult error) { + switch (error) { + case CUFFT_SUCCESS: + return "CUFFT_SUCCESS"; + + case CUFFT_INVALID_PLAN: + return "CUFFT_INVALID_PLAN"; + + case CUFFT_ALLOC_FAILED: + return "CUFFT_ALLOC_FAILED"; + + case CUFFT_INVALID_TYPE: + return "CUFFT_INVALID_TYPE"; + + case CUFFT_INVALID_VALUE: + return "CUFFT_INVALID_VALUE"; + + case CUFFT_INTERNAL_ERROR: + return "CUFFT_INTERNAL_ERROR"; + + case CUFFT_EXEC_FAILED: + return "CUFFT_EXEC_FAILED"; + + case CUFFT_SETUP_FAILED: + return "CUFFT_SETUP_FAILED"; + + case CUFFT_INVALID_SIZE: + return "CUFFT_INVALID_SIZE"; + + case CUFFT_UNALIGNED_DATA: + return "CUFFT_UNALIGNED_DATA"; + + case CUFFT_INCOMPLETE_PARAMETER_LIST: + return "CUFFT_INCOMPLETE_PARAMETER_LIST"; + + case CUFFT_INVALID_DEVICE: + return "CUFFT_INVALID_DEVICE"; + + case CUFFT_PARSE_ERROR: + return "CUFFT_PARSE_ERROR"; + + case CUFFT_NO_WORKSPACE: + return "CUFFT_NO_WORKSPACE"; + + case CUFFT_NOT_IMPLEMENTED: + return "CUFFT_NOT_IMPLEMENTED"; + + case CUFFT_LICENSE_ERROR: + return "CUFFT_LICENSE_ERROR"; + + case CUFFT_NOT_SUPPORTED: + return "CUFFT_NOT_SUPPORTED"; + } + + return ""; +} + +inline void CheckCuFFT(cufftResult err, const char* file, int line) { + if (err != CUFFT_SUCCESS) { + std::string error_msg = + std::string("cuFFT error = ") + CufftErrorStr(err) + " file = " + file + " line = " + std::to_string(line); + throw std::runtime_error(error_msg); + } +} + +inline void CheckCufftSize(size_t workspace_size, cufftHandle plan) { + size_t fft_workarea = 0; + auto cufft_err = cufftGetSize(plan, &fft_workarea); + if (cufft_err != CUFFT_SUCCESS || workspace_size < fft_workarea) { + throw std::runtime_error("workspace size not enough for FFT plan"); + } +} + +} diff --git a/cuda_util/cufft_plan.cpp b/cuda_util/cufft_plan.cc similarity index 99% rename from cuda_util/cufft_plan.cpp rename to cuda_util/cufft_plan.cc index 8a166c3..d077512 100644 --- a/cuda_util/cufft_plan.cpp +++ b/cuda_util/cufft_plan.cc @@ -9,7 +9,7 @@ */ #include "cufft_plan.h" -#include "checks.h" +#include "cufft_checks.h" namespace { int IntPow(int base, int power) { diff --git a/cuda_util/memory_policy.cc b/cuda_util/memory_policy.cc new file mode 100644 index 0000000..8083260 --- /dev/null +++ b/cuda_util/memory_policy.cc @@ -0,0 +1,42 @@ +/** +* ENVISAT and ERS ASAR instrument focusser for QA4EO activity (c) by CGI Estonia AS +* +* ENVISAT and ERS ASAR instrument focusser for QA4EO activity is licensed under a +* Creative Commons Attribution-ShareAlike 4.0 International License. +* +* You should have received a copy of the license along with this +* work. If not, see http://creativecommons.org/licenses/by-sa/4.0/ +*/ + +#include "memory_policy.h" + +namespace alus::cuda { + +MemoryAllocationForecast::MemoryAllocationForecast(size_t alignment) : alignment_{alignment} {} + +void MemoryAllocationForecast::Add(size_t bytes) { + if (bytes % alignment_ != 0) { + forecast_ += ((bytes / alignment_) + 1) * alignment_; + } else { + forecast_ += bytes; + } +} + +MemoryFitPolice::MemoryFitPolice(const CudaDevice& device, size_t percentage_allowed) + : device_{device}, percentage_{percentage_allowed}, total_memory_{device_.GetTotalGlobalMemory()} {} + +bool MemoryFitPolice::CanFit(size_t bytes) const { + const auto allowed_memory = + static_cast(total_memory_ * (static_cast(percentage_) / 100.0)); + if (bytes > allowed_memory) { + return false; + } + + if (device_.GetFreeGlobalMemory() < bytes) { + return false; + } + + return true; +} + +} // namespace alus::cuda \ No newline at end of file diff --git a/cuda_util/memory_policy.h b/cuda_util/memory_policy.h new file mode 100644 index 0000000..4406762 --- /dev/null +++ b/cuda_util/memory_policy.h @@ -0,0 +1,44 @@ +/** +* ENVISAT and ERS ASAR instrument focusser for QA4EO activity (c) by CGI Estonia AS +* +* ENVISAT and ERS ASAR instrument focusser for QA4EO activity is licensed under a +* Creative Commons Attribution-ShareAlike 4.0 International License. +* +* You should have received a copy of the license along with this +* work. If not, see http://creativecommons.org/licenses/by-sa/4.0/ +*/ +#pragma once + +#include + +#include "cuda_device.h" + +namespace alus::cuda { + +class MemoryAllocationForecast final { +public: + MemoryAllocationForecast() = delete; + explicit MemoryAllocationForecast(size_t alignment); + + void Add(size_t bytes); + [[nodiscard]] size_t Get() const { return forecast_; } + +private: + size_t alignment_; + size_t forecast_{}; +}; + +class MemoryFitPolice final { +public: + MemoryFitPolice() = delete; + MemoryFitPolice(const CudaDevice& device, size_t percentage_allowed); + + [[nodiscard]] bool CanFit(size_t bytes) const; + +private: + const CudaDevice& device_; + const size_t percentage_; + const size_t total_memory_; +}; + +} // namespace alus::cuda \ No newline at end of file diff --git a/envisat_format/src/ers_im_parse.cc b/envisat_format/src/ers_im_parse.cc index d83d2e4..bd16974 100644 --- a/envisat_format/src/ers_im_parse.cc +++ b/envisat_format/src/ers_im_parse.cc @@ -66,7 +66,8 @@ struct EchoMeta { // Auxiliary data and replica/calibration pulses uint64_t onboard_time; - uint16_t activity_task; + uint8_t activity_task; + uint8_t sample_flags; uint32_t image_format_counter; uint16_t pri_code; uint16_t swst_code; @@ -109,7 +110,7 @@ inline void FetchUint32(const uint8_t* array_start, uint32_t& var) { constexpr auto DR_NO_MAX{alus::asar::envformat::parseutil::MaxValueForBits()}; constexpr auto PACKET_COUNTER_MAX{alus::asar::envformat::parseutil::MaxValueForBits()}; constexpr uint8_t SUBCOMMUTATION_COUNTER_MAX{48}; -constexpr auto IMAGE_FORMAT_COUNTER_MAX{alus::asar::envformat::parseutil::MaxValueForBits()}; +constexpr auto IMAGE_FORMAT_COUNTER_MAX{alus::asar::envformat::parseutil::MaxValueForBits()}; void InitializeCounters(const uint8_t* packets_start, uint32_t& dr_no, uint8_t& packet_counter, uint8_t& subcommutation_counter, uint32_t& image_format_counter) { @@ -234,7 +235,7 @@ void ParseErsLevel0ImPackets(const std::vector& file_data, const DSD_lvl0& const auto end_filter_mjd = PtimeToMjd(packets_stop_filter); // https://earth.esa.int/eogateway/documents/20142/37627/ERS-products-specification-with-Envisat-format.pdf - // https://asf.alaska.edu/wp-content/uploads/2019/03/ers_ceos.pdf mixed between two. + // https://asf.alaska.edu/wp-content/uploads/2019/03/ers_ceos.pdf and ER-IS-ESA-GS-0002 mixed between three. for (size_t i = 0; i < mdsr.num_dsr; i++) { EchoMeta echo_meta = {}; #if HANDLE_DIAGNOSTICS @@ -286,28 +287,87 @@ void ParseErsLevel0ImPackets(const std::vector& file_data, const DSD_lvl0& } it = CopyBSwapPOD(echo_meta.data_record_number, it); + // ER-IS-ESA-GS-0002 4.4.2.4.5 + /* + * It is a binary counter which is incremented each time a new source packet of general IDHT header data is + * transmitted (about every 1 second). It is reset to zero with power-on of the IDHT-DCU. + * The first format transmitted after power-on will show the value "1". Bit O is defined as MSB. bit 7 as LSB. + */ echo_meta.packet_counter = it[0]; + /* + * binary counter counting the 8 byte segments from 1 to 48; it is reset for each new source packet. + * Bit o is defined as MSB, bit 7 as LSB. + */ echo_meta.subcommutation_counter = it[1]; it += 2; - it += 8; // IDHT General header source packet + /* + * These segments ot 8 bytes contain the subcommutated IDHT General Header source packet. + */ + it += 8; + // Table 4.4.2.6-2 in ER-IS-ESA-GS-0002 if (it[0] != 0xAA) { throw std::runtime_error(fmt::format( "ERS data packet's auxiliary section shall start with 0xAA, instead {:#x} was found for packet no {}", it[0], i + 1)); } it += 1; - it += 1; // OGRC/OBRC flag and Orbit ID code + // OBRC - On-Board Range Compressed | OGRC - On-Ground Range Compressed - See ER-IS-EPO-GS-0201 + it += 1; // bit 0 - OGRC/OBRC flag (0/1 respectively). Bits 1-4 Orbit number (0-15 -> 1-16) + /* + * The update of that binary counter shall occur every 4th PRI. + * The time relation to the echo data in the format is as following: + * the transfer of the ICU on-board time to the auxiliary memory + * occurs t2 before the RF transmit pulse as depicted in + * Fig. 4.4.2.4.6-3. The last significant bit is equal to 1/256 sec. + * + * ER-IS-ESA-GS-0002 - pg. 4.4 - 24 + */ echo_meta.onboard_time = 0; echo_meta.onboard_time |= static_cast(it[0]) << 24; echo_meta.onboard_time |= static_cast(it[1]) << 16; echo_meta.onboard_time |= static_cast(it[2]) << 8; echo_meta.onboard_time |= static_cast(it[3]) << 0; it += 4; - it = CopyBSwapPOD(echo_meta.activity_task, it); + /* + * This word shall define the activity task within the mode of + * operation, accompanied by the validity flag bits and the first + * sample flag bits. The update of the activity task word is controlled by the 4th PRI interrupt. + * ER-IS-ESA-GS-0002 - pg. 4.4 - 25 + * + * Pg. 4.4-25 + * Activity means generation of noise, echo, calibration or replica data. + * + * MSB LSB + * 10001000 - Noise; no calibration + * 10011001 - No echo; Cal. drift (EM only) + * 10101001 - Echo; Cal. drift + * 10101010 - Echo; Replica + * 10100000 - Echo; no Replica (because of OBRC) + */ + echo_meta.activity_task = it[0]; + /* + * ...continued from activity task's pg. 4.4-25 + * Bit + * 0 - echo invalid/valid (0/1) + * 1 - Cal. data/Replica data invalid/valid (0/1) + * 2 - Is noise + * 3 - Is Cal/Repl. + * 4 - Is echo + * ... spare bits + * + * The bits 2 to 4 (noise/cal/repl/echo) of byte 23 can only be set or reset with a 4 x PRI interval, + * therefore the start of a new activity is flagged for the first four formats. + */ + echo_meta.sample_flags = it[1]; + it += 2; + // Updated every format. Reset at the beginning of a transmission. it = CopyBSwapPOD(echo_meta.image_format_counter, it); it = CopyBSwapPOD(echo_meta.swst_code, it); it = CopyBSwapPOD(echo_meta.pri_code, it); + // Next item - Cal. S/S attenuation select - ER-IS-ESA-GS-0002 Table 4.4.2.4.6-2 + // Drift calibration data + // Lets skip all of that it += 194 + 10; const auto dr_no_gap = parseutil::CounterGap(last_dr_no, echo_meta.data_record_number); diff --git a/main.cc b/main.cc index 21f445e..584ed67 100644 --- a/main.cc +++ b/main.cc @@ -21,8 +21,9 @@ #include "alus_log.h" #include "args.h" #include "asar_constants.h" -#include "cuda_util/cufft_plan.h" -#include "cuda_util/device_padded_image.cuh" +#include "cufft_plan.h" +#include "device_padded_image.cuh" +#include "cuda_device_init.h" #include "envisat_aux_file.h" #include "envisat_lvl1_writer.h" #include "geo_tools.h" @@ -36,7 +37,6 @@ #include "sar/range_compression.cuh" #include "sar/range_doppler_algorithm.cuh" #include "sar/sar_chirp.h" -#include "status_assembly.h" namespace { struct IQ16 { @@ -78,6 +78,8 @@ int main(int argc, char* argv[]) { alus::asar::log::Initialize(); alus::asar::log::SetLevel(args.GetLogLevel()); + auto cuda_init = alus::cuda::CudaInit(); + LOGI << GetSoftwareVersion(); auto file_time_start = TimeStart(); const auto in_path{args.GetInputDsPath()}; @@ -118,6 +120,8 @@ int main(int argc, char* argv[]) { const auto target_product_type = alus::asar::specification::TryDetermineTargetProductFrom(product_type, args.GetFocussedProductType()); (void)target_product_type; + while(!cuda_init.IsFinished()); + cuda_init.CheckErrors(); alus::asar::envformat::ParseLevel0Packets(data, metadata, asar_meta, h_data, product_type, ins_file, asar_meta.sensing_start, asar_meta.sensing_stop); diff --git a/sar/range_compression.cu b/sar/range_compression.cu index 70b68d9..4ed9c6f 100755 --- a/sar/range_compression.cu +++ b/sar/range_compression.cu @@ -13,7 +13,7 @@ #include "cuda_util/cuda_cleanup.h" #include "cuda_util/cuda_util.h" #include "cuda_util/cufft_plan.h" -#include "checks.h" +#include "cufft_checks.h" __global__ void FrequencyDomainMultiply(cufftComplex* data_fft, const cufftComplex* chirp_fft, int range_fft_size, int azimuth_size) { @@ -80,7 +80,7 @@ void RangeCompression(DevicePaddedImage& data, const std::vector transpose + range FFT + transpose, 1 -> azimuth FFT @@ -298,7 +299,7 @@ void RangeDopplerAlgorithm(const SARMetadata& metadata, DevicePaddedImage& src_i src_img.Transpose(d_workspace); auto azimuth_fft = PlanRangeFFT(src_img.XStride(), src_img.YSize(), false); CufftPlanCleanup fft_cleanup(azimuth_fft); - CheckCufftSize(d_workspace.ByteSize(), azimuth_fft); + alus::cuda::CheckCufftSize(d_workspace.ByteSize(), azimuth_fft); CHECK_CUFFT_ERR(cufftSetWorkArea(azimuth_fft, d_workspace.Get())); CHECK_CUFFT_ERR(cufftExecC2C(azimuth_fft, src_img.Data(), src_img.Data(), CUFFT_FORWARD)); // TODO(priit) not actually needed, can fixed by inverting indexing in RCMC and Azimuth Ref multiply @@ -363,7 +364,7 @@ void RangeDopplerAlgorithm(const SARMetadata& metadata, DevicePaddedImage& src_i out_img.Transpose(d_workspace); auto azimuth_fft = PlanRangeFFT(out_img.XStride(), out_img.YSize(), false); CufftPlanCleanup fft_cleanup(azimuth_fft); - CheckCufftSize(d_workspace.ByteSize(), azimuth_fft); + alus::cuda::CheckCufftSize(d_workspace.ByteSize(), azimuth_fft); CHECK_CUFFT_ERR(cufftSetWorkArea(azimuth_fft, d_workspace.Get())); CHECK_CUFFT_ERR(cufftExecC2C(azimuth_fft, out_img.Data(), out_img.Data(), CUFFT_INVERSE)); out_img.Transpose(d_workspace); diff --git a/util/CMakeLists.txt b/util/CMakeLists.txt index f42a560..77734a8 100644 --- a/util/CMakeLists.txt +++ b/util/CMakeLists.txt @@ -23,6 +23,7 @@ target_link_libraries(util-static Boost::date_time Eigen3::Eigen app-static + cuda-util-static ) set_target_properties(util-static PROPERTIES diff --git a/util/include/checks.h b/util/include/checks.h index 27adaf6..c3888a3 100644 --- a/util/include/checks.h +++ b/util/include/checks.h @@ -9,8 +9,6 @@ */ #pragma once -#include - #include #include "alus_log.h" @@ -47,89 +45,6 @@ inline void CheckGdalError(CPLErr const err, char const* file, int const line) { } #endif -#define CHECK_CUFFT_ERR(x) CheckCuFFT(x, __FILE__, __LINE__) - -// copy paste from cuda-11.2/samples/common/inc/helper_cuda.h -inline const char* CufftErrorStr(cufftResult error) { - switch (error) { - case CUFFT_SUCCESS: - return "CUFFT_SUCCESS"; - - case CUFFT_INVALID_PLAN: - return "CUFFT_INVALID_PLAN"; - - case CUFFT_ALLOC_FAILED: - return "CUFFT_ALLOC_FAILED"; - - case CUFFT_INVALID_TYPE: - return "CUFFT_INVALID_TYPE"; - - case CUFFT_INVALID_VALUE: - return "CUFFT_INVALID_VALUE"; - - case CUFFT_INTERNAL_ERROR: - return "CUFFT_INTERNAL_ERROR"; - - case CUFFT_EXEC_FAILED: - return "CUFFT_EXEC_FAILED"; - - case CUFFT_SETUP_FAILED: - return "CUFFT_SETUP_FAILED"; - - case CUFFT_INVALID_SIZE: - return "CUFFT_INVALID_SIZE"; - - case CUFFT_UNALIGNED_DATA: - return "CUFFT_UNALIGNED_DATA"; - - case CUFFT_INCOMPLETE_PARAMETER_LIST: - return "CUFFT_INCOMPLETE_PARAMETER_LIST"; - - case CUFFT_INVALID_DEVICE: - return "CUFFT_INVALID_DEVICE"; - - case CUFFT_PARSE_ERROR: - return "CUFFT_PARSE_ERROR"; - - case CUFFT_NO_WORKSPACE: - return "CUFFT_NO_WORKSPACE"; - - case CUFFT_NOT_IMPLEMENTED: - return "CUFFT_NOT_IMPLEMENTED"; - - case CUFFT_LICENSE_ERROR: - return "CUFFT_LICENSE_ERROR"; - - case CUFFT_NOT_SUPPORTED: - return "CUFFT_NOT_SUPPORTED"; - } - - return ""; -} - -inline void CheckCuFFT(cufftResult err, const char* file, int line) { - if (err != CUFFT_SUCCESS) { - std::string error_msg = - std::string("cuFFT error = ") + CufftErrorStr(err) + " file = " + file + " line = " + std::to_string(line); - throw std::runtime_error(error_msg); - } -} - -inline void CheckNullptr(void* ptr, const char* file, int line) { - if (!ptr) { - std::string error_msg = std::string("nullptr file = ") + file + " line = " + std::to_string(line); - throw std::invalid_argument(error_msg); - } -} - -inline void CheckCufftSize(size_t workspace_size, cufftHandle plan) { - size_t fft_workarea = 0; - auto cufft_err = cufftGetSize(plan, &fft_workarea); - if (cufft_err != CUFFT_SUCCESS || workspace_size < fft_workarea) { - throw std::runtime_error("workspace size not enough for FFT plan"); - } -} - #define CHECK_BOOL(b) VerifyBool(b, __FILE__, __LINE__) inline void VerifyBool(bool cond, const char* file, int line) { diff --git a/util/include/img_output.h b/util/include/img_output.h index c02d50a..e999c4a 100644 --- a/util/include/img_output.h +++ b/util/include/img_output.h @@ -9,8 +9,6 @@ */ #pragma once -#include - -#include "cuda_util/device_padded_image.cuh" +#include "device_padded_image.cuh" void WriteIntensityPaddedImg(const DevicePaddedImage& img, const char* path); diff --git a/util/src/img_output.cc b/util/src/img_output.cc index 53d1358..9c41a23 100644 --- a/util/src/img_output.cc +++ b/util/src/img_output.cc @@ -10,18 +10,14 @@ #include "img_output.h" -#include #include #include #include -#include #include "alus_log.h" -#include "checks.h" -//#include "gdal_util.h" +#include "device_padded_image.cuh" #include "math_utils.h" -//#include "pugixml.hpp" namespace {