diff --git a/.docker/lyrical.arm64v8.dockerfile b/.docker/lyrical.arm64v8.dockerfile index 8d31b0a8..e96623b8 100644 --- a/.docker/lyrical.arm64v8.dockerfile +++ b/.docker/lyrical.arm64v8.dockerfile @@ -6,7 +6,7 @@ # setup below is the adapted version of what that base image contained, plus the xrdp # group-permission fix required by the verified Apple Silicon build. # - ROS_DISTRO=lyrical; ros-lyrical-ros-gz supplies the Gazebo Jetty vendor packages. -# - The current CUDA sonar remains disabled on ARM64; the open WGPU PR is intentionally excluded. +# - The sonar uses the portable WGPU backend when CUDA is unavailable on ARM64. # - ArduSub SITL includes the Python 3.14 compatibility shims validated on Ubuntu 26.04 # (imp/pipes modules, python-argparse removal, and PEP 668 handling). @@ -96,13 +96,15 @@ EXPOSE 22/tcp # --- ROS 2 Lyrical + Gazebo Jetty --- ARG ROS_DISTRO="lyrical" +ARG ROS_APT_SOURCE_VERSION="1.3.0" ENV GZ_VERSION=jetty RUN apt update && apt full-upgrade -y && apt autoremove -y # Gazebo Jetty is vendored by ros-lyrical-ros-gz on apt already — no separate Gazebo source build -RUN export ROS_APT_SOURCE_VERSION=$(curl -s https://api.github.com/repos/ros-infrastructure/ros-apt-source/releases/latest | grep -F "tag_name" | awk -F\" '{print $4}') && \ - curl -L -o /tmp/ros2-apt-source.deb \ +# Pin the bootstrap package to avoid unauthenticated GitHub API rate limits. +RUN curl --fail --show-error --location --retry 5 --retry-delay 2 \ + -o /tmp/ros2-apt-source.deb \ "https://github.com/ros-infrastructure/ros-apt-source/releases/download/${ROS_APT_SOURCE_VERSION}/ros2-apt-source_${ROS_APT_SOURCE_VERSION}.$(. /etc/os-release && echo ${UBUNTU_CODENAME:-${VERSION_CODENAME}})_all.deb" && \ dpkg -i /tmp/ros2-apt-source.deb && \ apt update && \ diff --git a/README.md b/README.md index 9c3c17c1..04cebfe2 100644 --- a/README.md +++ b/README.md @@ -19,6 +19,14 @@ Start with [installation](https://ioes-lab.github.io/POSIM/install.html) or [Docker](https://ioes-lab.github.io/POSIM/docker.html), then run your [first simulation](https://ioes-lab.github.io/POSIM/quickstart.html). +## Multibeam sonar + +Run the BlueView P900 sensor with WGPU, CUDA or CPU processing. Select the backend +at launch and inspect sonar images, raw acoustic data and point clouds. + +- English: [sensor and demo](website/content/en/sonar.md), [build and backend setup](website/content/en/sonar-tuning.md), [Apple Silicon](website/content/en/sonar-tuning.md#apple-silicon) +- 한국어: WGPU·CUDA·CPU 계산 방식을 선택하여 소나 영상·원시 음향 데이터·점군을 확인합니다. [센서와 예제](website/content/ko/sonar.md), [빌드와 계산 설정](website/content/ko/sonar-tuning.md), [Apple Silicon](website/content/ko/sonar-tuning.md#apple-silicon) + ## Try WWW-POSIM [WWW-POSIM](https://www-posim.vercel.app/) uses POSIM as its simulation engine. diff --git a/docs/engineering/sonar-public-documentation.md b/docs/engineering/sonar-public-documentation.md new file mode 100644 index 00000000..93a2bf7a --- /dev/null +++ b/docs/engineering/sonar-public-documentation.md @@ -0,0 +1,40 @@ +# Sonar public documentation evidence + +## Source snapshot + +- POSIM sonar PR: https://github.com/IOES-Lab/POSIM/pull/6 +- Source head: `ad6165bd6c2b4e407f18ecb264e20bb76da8d4ca`. +- Input guides: `gazebo/DEMO_GUIDE.md` and `gazebo/DEMO_GUIDE_AppleSilicon_MacOSX.md`. +- Runtime report: https://github.com/IOES-Lab/POSIM/pull/6#issuecomment-6135206430. +- User clarified that these documentation changes belong in PR #6; PR #7 is the closed ocean-waves contribution. + +## Source-backed user behavior + +- `sonar_compute_cpu.cc`: `auto` tries CUDA, WGPU, CPU; explicit failures and invalid names return no backend. +- `MultibeamSonarSensor.cc`: environment name is `POSIM_SONAR_COMPUTE_BACKEND`; publishes Image and ProjectedSonarImage directly to ROS 2, and PointCloudPacked over Gazebo transport; rendering is required for GpuRays. +- Demo launch: declares `compute_backend` (default `auto`) and `rviz`, and sets the POSIM backend environment variable. +- Sonar CMake: WGPU is required, CUDA is optional, `CUDA_ARCHITECTURE` defaults to `native` and forwards to `CMAKE_CUDA_ARCHITECTURES`; Apple builds link Metal/QuartzCore. +- Rust pipeline: Vulkan is the Linux default, Metal is the macOS default; supports adapter substring, backend API and power-preference selection; GPU context is shared in-process. +- WGPU 0.20.1 power-preference values checked against https://github.com/gfx-rs/wgpu/blob/v0.20.1/wgpu/src/util/init.rs: `high`, `low`, `none`. + +## Documentation decisions + +Public pages describe present behavior and commands. The two demo files and sonar README point to the authoritative English/Korean guides. Root README and installation, requirements, overview and troubleshooting pages use the same supported backend facts. + +The inherited demo guides refer to Ubuntu 24.04/ROS Rolling and macOS Jazzy/Harmonic setup. These version-specific setup commands are not copied into Lyrical/Jetty instructions. Native macOS instructions state the matching ROS/Gazebo/rendering prerequisites; they are not a turnkey host installer or a fresh native runtime validation. + +The contributor reports Ubuntu 26.04.1, ROS 2 Lyrical, Gazebo 10.5.0, RTX 4060 and Intel iGPU runtime tests. Reported panel range errors at 2–8 m and 0–45 degrees were approximately 0.03–0.05 m for WGPU and 0.03–0.06 m for CUDA, with 2.5 cm range bins. These are attributed test observations, not new measurements by this documentation change or public accuracy guarantees. + +The public text omits migration history, commit counts, implementation cleanup, shared-buffer allocation counts and before/after comparisons. Build/check and browser review results are recorded separately from user instructions. + +## Documentation validation + +- Node.js 24: `pnpm catalog`, `pnpm build`, `pnpm check`. +- Full repository check: 50 pages, 3,575 local links/anchors, 38 bilingual public media, bilingual search and 42 source launch references. +- Website-only build/check passes without the ROS/model source tree. +- EN/KO command blocks match for sonar, sonar tuning, installation, requirements, overview and troubleshooting. Sonar Bash blocks pass `bash -n`. +- Both demo wrappers pass the public-wording checker. `git diff --check` passes. +- WebKit desktop (1440 x 1000) and mobile (390 x 844) checks cover both sonar pages in both languages: no document overflow or JavaScript page errors, and language switching passes. +- Desktop Apple Silicon section and Korean mobile demo output visually inspected. Browser captures/results are local ignored files under `website/.validation/sonar/`. +- Website documentation-basis metadata uses the inspected PR source snapshot; no Notion review date or source media metadata is changed. +- These are documentation checks. This change does not claim a new sonar simulation run or a native Apple Silicon build. diff --git a/extras/ros-lyrical-gz-jetty-install.sh b/extras/ros-lyrical-gz-jetty-install.sh index b199364d..60a69229 100755 --- a/extras/ros-lyrical-gz-jetty-install.sh +++ b/extras/ros-lyrical-gz-jetty-install.sh @@ -59,8 +59,10 @@ sudo apt update && sudo apt install -y jq # shellcheck disable=SC1091 . /etc/os-release UBUNTU_CODENAME="${UBUNTU_CODENAME:-${VERSION_CODENAME}}" -ROS_APT_SOURCE_VERSION=$(curl -s https://api.github.com/repos/ros-infrastructure/ros-apt-source/releases/latest | jq -r '.tag_name') && \ - curl -L -o /tmp/ros2-apt-source.deb \ +# Use a known release instead of a rate-limited unauthenticated API lookup. +ROS_APT_SOURCE_VERSION="${ROS_APT_SOURCE_VERSION:-1.3.0}" +curl --fail --show-error --location --retry 5 --retry-delay 2 \ + -o /tmp/ros2-apt-source.deb \ "https://github.com/ros-infrastructure/ros-apt-source/releases/download/${ROS_APT_SOURCE_VERSION}/ros2-apt-source_${ROS_APT_SOURCE_VERSION}.${UBUNTU_CODENAME}_all.deb" && \ sudo apt-get install -y /tmp/ros2-apt-source.deb && \ rm -f /tmp/ros2-apt-source.deb diff --git a/gazebo/DEMO_GUIDE.md b/gazebo/DEMO_GUIDE.md new file mode 100644 index 00000000..f72c1682 --- /dev/null +++ b/gazebo/DEMO_GUIDE.md @@ -0,0 +1,12 @@ +# Multibeam sonar demo + +Run the BlueView P900 demo with WGPU, CUDA or CPU processing. The public guides +cover build prerequisites, launch commands, backend selection, output topics +and troubleshooting. + +- English: [run the demo](../website/content/en/sonar.md), [build and backend setup](../website/content/en/sonar-tuning.md) +- 한국어: [소나 예제 실행](../website/content/ko/sonar.md), [빌드와 계산 설정](../website/content/ko/sonar-tuning.md) + +POSIM uses ROS 2 Lyrical and Gazebo Jetty. Start with the +[installation guide](../website/content/en/install.md) or +[한국어 설치 안내](../website/content/ko/install.md), then follow the sonar guides above. diff --git a/gazebo/DEMO_GUIDE_AppleSilicon_MacOSX.md b/gazebo/DEMO_GUIDE_AppleSilicon_MacOSX.md new file mode 100644 index 00000000..4b342df4 --- /dev/null +++ b/gazebo/DEMO_GUIDE_AppleSilicon_MacOSX.md @@ -0,0 +1,11 @@ +# Multibeam sonar on Apple Silicon + +The public sonar build guide covers Metal for native macOS and CPU acoustic +processing in a Linux ARM64 container. Both require Gazebo rendering and a +matching ROS 2 / Gazebo environment. + +- English: [Apple Silicon setup](../website/content/en/sonar-tuning.md#apple-silicon), [sensor and demo](../website/content/en/sonar.md) +- 한국어: [Apple Silicon 설정](../website/content/ko/sonar-tuning.md#apple-silicon), [센서와 예제](../website/content/ko/sonar.md) + +Use the [Docker guide](../website/content/en/docker.md) or +[한국어 Docker 안내](../website/content/ko/docker.md) for a Linux ARM64 environment. diff --git a/gazebo/posim_gz_multibeam_sonar/README.md b/gazebo/posim_gz_multibeam_sonar/README.md index 1fabca55..ca7839ab 100644 --- a/gazebo/posim_gz_multibeam_sonar/README.md +++ b/gazebo/posim_gz_multibeam_sonar/README.md @@ -1,20 +1,25 @@ # POSIM multibeam sonar -The ray-based point-scattering model produces intensity/range data, with -phase, reverberation and speckle effects. +The ray-based point-scattering model produces intensity–range data, sonar images +and point clouds. Select WGPU, CUDA or CPU processing at launch. WGPU uses Vulkan +on Linux and Metal on native macOS; all backends require Gazebo rendering. -## Backend requirements +## Documentation -- ROS 2 Lyrical and Gazebo Jetty -- Compatible NVIDIA GPU and CUDA toolkit -- Built CUDA sonar libraries +- [Sensor, demo and output topics](../../website/content/en/sonar.md) +- [Build, backend selection and troubleshooting](../../website/content/en/sonar-tuning.md) +- [Apple Silicon](../../website/content/en/sonar-tuning.md#apple-silicon) +- [POSIM installation](../../website/content/en/install.md) +- [Research citation and licenses](../../website/content/en/citation.md) -CMake skips CUDA targets when the toolkit is unavailable. Check library -installation and sensor output with the [sonar build guide](../../website/content/en/sonar-tuning.md). +## 한국어 안내 -## Documentation +광선 기반 점 산란 모델로 강도·거리 데이터, 소나 영상과 점군을 생성합니다. +실행할 때 WGPU·CUDA·CPU 계산 방식을 선택합니다. WGPU는 Linux에서 Vulkan, +macOS에서 직접 실행할 때 Metal을 사용합니다. 모든 방식에서 Gazebo 렌더링이 필요합니다. -See the [POSIM sonar guide](../../website/content/en/sonar.md) -for the model and configuration. Use the -[POSIM installation guide](../../website/content/en/install.md) for current repository -and workspace setup. +- [센서·예제·출력 토픽](../../website/content/ko/sonar.md) +- [빌드·계산 방식·문제 해결](../../website/content/ko/sonar-tuning.md) +- [Apple Silicon](../../website/content/ko/sonar-tuning.md#apple-silicon) +- [POSIM 설치](../../website/content/ko/install.md) +- [연구 인용과 라이선스](../../website/content/ko/citation.md) diff --git a/gazebo/posim_gz_multibeam_sonar/multibeam_sonar/CMakeLists.txt b/gazebo/posim_gz_multibeam_sonar/multibeam_sonar/CMakeLists.txt index 1f01ac18..737c637d 100644 --- a/gazebo/posim_gz_multibeam_sonar/multibeam_sonar/CMakeLists.txt +++ b/gazebo/posim_gz_multibeam_sonar/multibeam_sonar/CMakeLists.txt @@ -1,74 +1,105 @@ cmake_minimum_required(VERSION 3.11.0 FATAL_ERROR) # Suppress developer warnings for the entire workspace. - set(CMAKE_SUPPRESS_DEVELOPER_WARNINGS TRUE CACHE BOOL "Suppress developer warnings" FORCE) -# Set CMP0144 to NEW to ensure find_package -# uses upper-case _ROOT variables. +# Ensure find_package uses upper-case _ROOT variables. cmake_policy(SET CMP0144 NEW) -project(multibeam_sonar) +set(CUDA_ARCHITECTURE "native" CACHE STRING + "Target CUDA SM version(s), e.g. native or 70;80;89") +set(CMAKE_CUDA_ARCHITECTURES "${CUDA_ARCHITECTURE}") -set(CUDA_ARCHITECTURE "60" CACHE STRING "Target CUDA SM version") +project(multibeam_sonar) find_package(ament_cmake REQUIRED) +find_package(gz-cmake 5 REQUIRED) +find_package(gz-sim 10 REQUIRED) +find_package(gz-sensors 10 REQUIRED) +find_package(gz-rendering 10 REQUIRED OPTIONAL_COMPONENTS ogre ogre2) +find_package(gz-transport 15 REQUIRED) +find_package(gz-msgs 12 REQUIRED) +find_package(rclcpp REQUIRED) +find_package(sensor_msgs REQUIRED) +find_package(geometry_msgs REQUIRED) +find_package(OpenCV REQUIRED) +find_package(marine_acoustic_msgs REQUIRED) +find_package(Eigen3 REQUIRED) +find_package(wgpu_vendor REQUIRED) find_package(CUDAToolkit QUIET) if(CUDAToolkit_FOUND) - + set(CMAKE_CUDA_COMPILER "${CUDAToolkit_NVCC_EXECUTABLE}" + CACHE FILEPATH "CUDA compiler" FORCE) enable_language(CUDA) find_package(CUDA REQUIRED) - message(STATUS "CUDA found, enabling CUDA support.") + message(STATUS + "CUDA ${CUDAToolkit_VERSION} found (SM ${CUDA_ARCHITECTURE}), " + "enabling CUDA support.") include_directories(${CUDA_INCLUDE_DIRS}) - set(CUDA_NVCC_FLAGS "${CUDA_NVCC_FLAGS} -arch=sm_${CUDA_ARCHITECTURE}") - find_package(tinyxml2 CONFIG REQUIRED) - find_package(gz-cmake 5 REQUIRED) - find_package(gz-sim 10 REQUIRED) - find_package(gz-sensors 10 REQUIRED) - find_package(gz-rendering 10 REQUIRED OPTIONAL_COMPONENTS ogre ogre2) - find_package(gz-transport 15 REQUIRED) - find_package(gz-msgs 12 REQUIRED) - find_package(rclcpp REQUIRED) - find_package(sensor_msgs REQUIRED) - find_package(geometry_msgs REQUIRED) - find_package(rosidl_default_generators REQUIRED) - find_package(OpenCV REQUIRED) - find_package(marine_acoustic_msgs REQUIRED) - find_package(cv_bridge REQUIRED) - - if(TARGET gz-rendering::ogre) - set(HAVE_OGRE TRUE) - add_definitions(-DWITH_OGRE) - endif() +endif() - if(TARGET gz-rendering::ogre2) - set(HAVE_OGRE2 TRUE) - add_definitions(-DWITH_OGRE2) - endif() +if(TARGET gz-rendering::ogre) + add_definitions(-DWITH_OGRE) +endif() - add_library(${PROJECT_NAME} - SHARED - MultibeamSonarSensor.cc - sonar_calculation_cuda.cu - ) +if(TARGET gz-rendering::ogre2) + add_definitions(-DWITH_OGRE2) +endif() + +set(SONAR_SOURCES + MultibeamSonarSensor.cc + sonar_compute_cpu.cc + sonar_compute_wgpu.cc + sonar_compute_cuda.cc +) + +if(CUDAToolkit_FOUND) + list(APPEND SONAR_SOURCES sonar_calculation_cuda.cu) +endif() +add_library(${PROJECT_NAME} SHARED ${SONAR_SOURCES}) + +if(CUDAToolkit_FOUND) set_target_properties(${PROJECT_NAME} PROPERTIES - CUDA_SEPARABLE_COMPILATION ON + CUDA_SEPARABLE_COMPILATION ON ) + target_compile_definitions(${PROJECT_NAME} PRIVATE POSIM_HAS_CUDA_BACKEND=1) +endif() - target_include_directories(${PROJECT_NAME} - PUBLIC - ${gz-sensors_INCLUDE_DIRS} - ${gz-rendering_INCLUDE_DIRS} - ${gz-msgs_INCLUDE_DIRS} - ${CUDA_INCLUDE_DIRS} - ${gz-transport_INCLUDE_DIRS} - ${gz-sim_INCLUDE_DIRS} +target_include_directories(${PROJECT_NAME} PUBLIC + ${rclcpp_INCLUDE_DIRS} + ${sensor_msgs_INCLUDE_DIRS} + ${geometry_msgs_INCLUDE_DIRS} + ${OpenCV_INCLUDE_DIRS} + ${EIGEN3_INCLUDE_DIRS} + ${marine_acoustic_msgs_INCLUDE_DIRS} +) + +target_link_libraries(${PROJECT_NAME} + gz-sim::gz-sim + gz-sensors::gz-sensors + gz-msgs::gz-msgs + gz-transport::gz-transport + gz-rendering::gz-rendering + ${rclcpp_LIBRARIES} + ${sensor_msgs_LIBRARIES} + ${geometry_msgs_LIBRARIES} + ${OpenCV_LIBS} + Eigen3::Eigen + ${marine_acoustic_msgs_LIBRARIES} +) + +if(APPLE) + target_link_libraries(${PROJECT_NAME} + "-framework Metal" + "-framework QuartzCore" ) +endif() +if(CUDAToolkit_FOUND) find_library(CUBLAS_LIB cublas HINTS ${CUDAToolkit_LIBRARY_DIR} @@ -79,43 +110,44 @@ if(CUDAToolkit_FOUND) if(NOT CUBLAS_LIB) message(FATAL_ERROR "cuBLAS not found") endif() - target_link_libraries(${PROJECT_NAME} - gz-sim::gz-sim - gz-sensors::gz-sensors - gz-msgs::gz-msgs - gz-transport::gz-transport - gz-rendering::gz-rendering ${CUDA_LIBRARIES} ${CUDA_CUFFT_LIBRARIES} ${CUBLAS_LIB} ) +endif() - if(TARGET gz-sensors::gz-sensors-rendering) - target_link_libraries(${PROJECT_NAME} gz-sensors::gz-sensors-rendering) - endif() +set(_wgpu_target_linked FALSE) +if(TARGET wgpu::wgpu_vendor) + target_link_libraries(${PROJECT_NAME} wgpu::wgpu_vendor) + set(_wgpu_target_linked TRUE) +elseif(TARGET wgpu_vendor) + target_link_libraries(${PROJECT_NAME} wgpu_vendor) + set(_wgpu_target_linked TRUE) +endif() - install(TARGETS ${PROJECT_NAME} - DESTINATION lib/${PROJECT_NAME} - ) +# Ensure the Rust static archive is linked into this shared library. The +# exported interface target may not always propagate non-installed imported +# targets across package boundaries in overlay builds. +if(NOT _wgpu_target_linked AND DEFINED wgpu_vendor_DIR) + get_filename_component( + _wgpu_vendor_prefix "${wgpu_vendor_DIR}/../../.." ABSOLUTE) + set(_sonar_wgpu_static "${_wgpu_vendor_prefix}/lib/libsonar_wgpu.a") + if(EXISTS "${_sonar_wgpu_static}") + target_link_libraries(${PROJECT_NAME} "${_sonar_wgpu_static}") + endif() +endif() - ament_target_dependencies(${PROJECT_NAME} - rclcpp - sensor_msgs - rosidl_default_generators - OpenCV - marine_acoustic_msgs - geometry_msgs - cv_bridge - ) +if(TARGET gz-sensors::gz-sensors-rendering) + target_link_libraries(${PROJECT_NAME} gz-sensors::gz-sensors-rendering) +endif() - ament_environment_hooks( - "${CMAKE_CURRENT_SOURCE_DIR}/hooks/${PROJECT_NAME}.dsv.in" - ) +install(TARGETS ${PROJECT_NAME} + DESTINATION lib/${PROJECT_NAME} +) -else() - message(STATUS "CUDA Toolkit not found or disabled: " - "Skipping CUDA-specific targets") -endif() +ament_environment_hooks( + "${CMAKE_CURRENT_SOURCE_DIR}/hooks/${PROJECT_NAME}.dsv.in" +) ament_package() diff --git a/gazebo/posim_gz_multibeam_sonar/multibeam_sonar/MultibeamSonarSensor.cc b/gazebo/posim_gz_multibeam_sonar/multibeam_sonar/MultibeamSonarSensor.cc index 07ae98db..dd63cf1a 100644 --- a/gazebo/posim_gz_multibeam_sonar/multibeam_sonar/MultibeamSonarSensor.cc +++ b/gazebo/posim_gz_multibeam_sonar/multibeam_sonar/MultibeamSonarSensor.cc @@ -46,15 +46,16 @@ #include #include #include "MultibeamSonarSensor.hh" -#include "sonar_calculation_cuda.cuh" #include -#include +#include +#include #include #include #include #include #include +#include namespace gz { @@ -142,7 +143,7 @@ class InMemoryTimeVaryingVectorField } if (this->zData && this->zSession) { - const auto interpolation = this->zData->LookUp(this->ySession.value(), _pos); + const auto interpolation = this->zData->LookUp(this->zSession.value(), _pos); outcome.Z(interpolation.value_or(0.)); } return outcome; @@ -173,6 +174,41 @@ class InMemoryTimeVaryingVectorField const GridT * zData{nullptr}; }; +/// \brief Copy an 8-bit OpenCV image into a ROS image message. +/// +/// The sonar only publishes BGR8 and RGB8 visualization images here. Keeping +/// this small conversion local avoids requiring cv_bridge for a byte-for-byte +/// copy and prevents OpenCV major-version coupling at runtime. +void CopyImageToMessage( + const cv::Mat & _image, const std::string & _encoding, sensor_msgs::msg::Image & _message) +{ + _message.height = static_cast(_image.rows); + _message.width = static_cast(_image.cols); + _message.encoding = _encoding; + _message.is_bigendian = false; + _message.step = static_cast(_image.cols * _image.elemSize()); + + const auto dataSize = static_cast(_message.step) * _message.height; + _message.data.resize(dataSize); + if (dataSize == 0) + { + return; + } + + if (_image.isContinuous()) + { + std::memcpy(_message.data.data(), _image.data, dataSize); + return; + } + + for (int row = 0; row < _image.rows; ++row) + { + std::memcpy( + _message.data.data() + static_cast(row) * _message.step, _image.ptr(row), + _message.step); + } +} + } // namespace using namespace gz::msgs; @@ -183,6 +219,21 @@ MultibeamSonarSensor::MultibeamSonarSensor() : dataPtr(new Implementation()) {} ////////////////////////////////////////////////// MultibeamSonarSensor::~MultibeamSonarSensor() { + // Stop the background compute thread before tearing down sensors. + { + std::lock_guard lk(this->dataPtr->computeMutex_); + this->dataPtr->stopThread_ = true; + } + this->dataPtr->computeCV_.notify_one(); + if (this->dataPtr->computeThread_.joinable()) + { + this->dataPtr->computeThread_.join(); + } + + if (this->dataPtr->ros_executor_ && this->dataPtr->ros_node_) + { + this->dataPtr->ros_executor_->remove_node(this->dataPtr->ros_node_); + } this->dataPtr->rayConnection.reset(); this->dataPtr->sceneChangeConnection.reset(); // CSV log write stream close @@ -278,6 +329,15 @@ bool MultibeamSonarSensor::Implementation::Initialize(MultibeamSonarSensor * _se { gzerr << "Failed to initialize beam arrangement for " << "[" << _sensor->Name() << "] sensor." << std::endl; + // An unconfigured GpuRays left in the scene makes Ogre throw on the next render + // and aborts the server, so remove it and fail only this sensor. + this->rayConnection.reset(); + if (this->raySensor) + { + _sensor->Scene()->DestroySensor(this->raySensor); + this->raySensor.reset(); + } + this->computeBackend.reset(); return false; } @@ -286,6 +346,24 @@ bool MultibeamSonarSensor::Implementation::Initialize(MultibeamSonarSensor * _se this->referenceFrameRotation = referenceFrameTransform.Rot().Inverse(); + // Spawn background sonar compute thread now that all setup is complete. + this->computeThread_ = std::thread( + [this]() + { + while (true) + { + std::unique_lock lk(this->computeMutex_); + this->computeCV_.wait(lk, [this] { return this->newFrameReady_ || this->stopThread_; }); + if (this->stopThread_) + { + break; + } + this->newFrameReady_ = false; + lk.unlock(); + this->ComputeSonarImage(); + } + }); + gzmsg << "Initialized [" << _sensor->Name() << "] sensor." << std::endl; this->initialized = true; return true; @@ -297,6 +375,28 @@ bool MultibeamSonarSensor::Implementation::InitializeBeamArrangement(MultibeamSo this->beams.clear(); this->beamTargets.clear(); + // Select the compute backend before adding anything to the scene, so an unavailable + // backend fails this sensor without leaving a half-configured GpuRays behind. + if (const char * backendEnv = std::getenv("POSIM_SONAR_COMPUTE_BACKEND")) + { + this->requestedBackend = backendEnv; + gzmsg << "POSIM_SONAR_COMPUTE_BACKEND=" << this->requestedBackend << std::endl; + } + else + { + this->requestedBackend = "auto"; + gzmsg << "POSIM_SONAR_COMPUTE_BACKEND not set. Using default backend=auto" << std::endl; + } + + this->computeBackend = CreateComputeBackend(this->requestedBackend); + if (!this->computeBackend) + { + gzerr << "Unable to initialize requested sonar backend [" << this->requestedBackend << "]." + << std::endl; + return false; + } + gzmsg << "Using sonar compute backend: " << this->computeBackend->Name() << std::endl; + this->raySensor = _sensor->Scene()->CreateGpuRays(_sensor->Name() + "_ray_sensor"); if (!this->raySensor) { @@ -383,19 +483,14 @@ bool MultibeamSonarSensor::Implementation::InitializeBeamArrangement(MultibeamSo sensorElement->Get("sonarImageTopicName", "sonar_image").first; gzmsg << "sonarImageTopicName: " << this->sonarImageTopicName << std::endl; + // NOTE: frameName is the TF frame name for published sonar messages (frame_id in ROS headers). + // This is explicitly configured in SDF rather than derived from Gazebo's + // internal sensor frame ID. Must match TF tree for RViz transforms to work correctly. + // Verify in SDF files that matches the expected TF frame name. this->frameName = sensorElement->Get("frameName", "forward_sonar_optical_link").first; gzmsg << "frameName: " << this->frameName << std::endl; - this->frameId = _sensor->FrameId(); - size_t pos = 0; - while ((pos = this->frameId.find("::", pos)) != std::string::npos) - { - this->frameId.replace(pos, 2, "/"); - pos += 1; - } - gzmsg << "frameId: " << this->frameId << std::endl; - // ROS Initialization if (!rclcpp::ok()) @@ -404,7 +499,8 @@ bool MultibeamSonarSensor::Implementation::InitializeBeamArrangement(MultibeamSo } this->ros_node_ = std::make_shared("multibeam_sonar_node"); - this->ros_node_->set_parameters({rclcpp::Parameter("use_sim_time", true)}); + this->ros_executor_ = std::make_shared(); + this->ros_executor_->add_node(this->ros_node_); // Create the point cloud publisher this->pointPub = this->node.Advertise( @@ -689,12 +785,47 @@ bool MultibeamSonarSensor::Implementation::InitializeBeamArrangement(MultibeamSo this->constMu = true; this->mu = 1e-3; + SonarComputeInput prototype; + prototype.depthImage = &this->pointCloudImage; + prototype.nBeams = this->nBeams; + prototype.nRays = this->nRays; + prototype.nFreq = this->nFreq; + prototype.raySkips = this->raySkips; + prototype.hFOV = this->hFOV; + prototype.vFOV = this->vFOV; + prototype.maxDistance = this->maxDistance; + prototype.soundSpeed = this->soundSpeed; + prototype.attenuation = this->attenuation; + prototype.sensorGain = this->sensorGain; + prototype.debugFlag = this->debugFlag; + prototype.blazingFlag = this->blazingFlag; + prototype.window = this->window; + prototype.rangeVector = this->rangeVector; + prototype.beamCorrector = this->beamCorrector; + prototype.beamCorrectorSum = this->beamCorrectorSum; + prototype.sourceLevel = this->sourceLevel; + prototype.bandwidth = this->bandwidth; + prototype.sonarFreq = this->sonarFreq; + prototype.elevation_angles.assign(this->elevation_angles, this->elevation_angles + this->nRays); + prototype.seed = this->sonarSeed; + + if (!this->computeBackend->Initialize(prototype)) + { + gzerr << "Failed to initialize sonar backend [" << this->computeBackend->Name() << "]." + << std::endl; + return false; + } + return true; } ////////////////////////////////////////////////// std::vector MultibeamSonarSensor::RenderingSensors() const { + if (!this->dataPtr->raySensor) + { + return {}; + } return {this->dataPtr->raySensor}; } @@ -724,9 +855,14 @@ void MultibeamSonarSensor::Implementation::OnNewFrame( // Fill point cloud with the ray buffer this->FillPointCloudMsg(this->rayBuffer); + // Signals the background compute thread to process the new frame if (this->pointCloudImage.size().width != 0) { - this->ComputeSonarImage(); + { + std::lock_guard lk(this->computeMutex_); + this->newFrameReady_ = true; + } + this->computeCV_.notify_one(); } } @@ -805,7 +941,10 @@ void MultibeamSonarSensor::PostUpdate(const std::chrono::steady_clock::duration << "cannot estimate velocities." << std::endl; return; } - rclcpp::spin_some(this->dataPtr->ros_node_); + if (this->dataPtr->ros_executor_) + { + this->dataPtr->ros_executor_->spin_some(); + } if (this->dataPtr->publishingPointCloud) { @@ -998,9 +1137,11 @@ cv::Mat MultibeamSonarSensor::Implementation::ComputeNormalImage(cv::Mat & depth images.at(0) = n1; // for green channel images.at(1) = n2; // for red channel - // Calculate focal length (?) Not sure if this is the right way; - // I could not find in the original DAVE the definition of the focal length - double focal_length = (0.5 * this->pointMsg.width()) / tan(0.5 * this->hFOV); + // Calculate focal length using the actual depth image dimensions (cols) so + // this function is safe to call from a background thread without accessing + // the shared pointMsg field. Avoid data race with pointMsg.width() which may + // be being written to by the render thread concurrently. + double focal_length = (0.5 * depth.cols) / tan(0.5 * this->hFOV); images.at(2) = 1.0 / focal_length * depth; // for blue channel cv::Mat normal_image; @@ -1020,14 +1161,19 @@ cv::Mat MultibeamSonarSensor::Implementation::ComputeNormalImage(cv::Mat & depth } // Precalculation of corrector sonar calculation -void MultibeamSonarSensor::Implementation::ComputeCorrector() +void MultibeamSonarSensor::Implementation::ComputeCorrector(int _snapshotWidth, int _nBeams) { - double hPixelSize = this->hFOV / (this->pointMsg.width() - 1); + if (_snapshotWidth <= 1 || _nBeams <= 0) + { + return; + } + + double hPixelSize = this->hFOV / (_snapshotWidth - 1); // Beam culling correction precalculation - for (size_t beam = 0; beam < this->nBeams; beam++) + for (int beam = 0; beam < _nBeams; beam++) { - for (size_t beam_other = 0; beam_other < this->nBeams; beam_other++) + for (int beam_other = 0; beam_other < _nBeams; beam_other++) { float azimuthBeamPattern = unnormalized_sinc( M_PI * 0.884 / hPixelSize * @@ -1043,21 +1189,51 @@ void MultibeamSonarSensor::Implementation::ComputeCorrector() void MultibeamSonarSensor::Implementation::ComputeSonarImage() { - this->lock_.lock(); - cv::Mat normal_image = this->ComputeNormalImage(this->pointCloudImage); - double vPixelSize = this->vFOV / (this->pointMsg.height() - 1); - double hPixelSize = this->hFOV / (this->pointMsg.width() - 1); - - if (this->beamCorrectorSum == 0) + // Snapshot the depth image quickly, then release the lock so the render + // thread can fill the next frame without waiting for GPU compute. + cv::Mat depthSnapshot; + // Capture the timestamp and image dimensions while holding the lock so that + // all published messages share one coherent stamp matching when the data was + // recorded -- mismatched stamps are the primary cause of RViz flicker. + rclcpp::Time capturedStamp; { - ComputeCorrector(); - } + std::lock_guard snapshot_lk(this->lock_); + if (this->pointCloudImage.empty()) + { + return; + } + depthSnapshot = this->pointCloudImage.clone(); + capturedStamp = this->ros_node_->now(); - if (this->reflectivityImage.rows == 0) - { - this->reflectivityImage = - cv::Mat(this->pointMsg.width(), this->pointMsg.height(), CV_32FC1, cv::Scalar(this->mu)); + if (this->beamCorrectorSum == 0) + { + ComputeCorrector(depthSnapshot.cols, this->nBeams); + } + + if (this->reflectivityImage.rows == 0) + { + this->reflectivityImage = + cv::Mat(depthSnapshot.rows, depthSnapshot.cols, CV_32FC1, cv::Scalar(this->mu)); + } } + // lock_ released -> render thread can proceed immediately. + + // NOTE: reflectivityImage is used here without a lock. + // This is safe for now because only the compute thread writes to it + // and the render thread only reads it. If the render thread ever starts modifying + // reflectivityImage, this will cause a data race. + // TODO: Either guarantee that reflectivityImage stays read-only + // after initialization, or add proper locking around the compute step. + + // Use snapshot dimensions from here on to avoid racing with the render thread + // that may be writing pointMsg concurrently. + const int snapWidth = depthSnapshot.cols; + const int snapHeight = depthSnapshot.rows; + + double vPixelSize = this->vFOV / (snapHeight - 1); + double hPixelSize = this->hFOV / (snapWidth - 1); + + cv::Mat normal_image = this->ComputeNormalImage(depthSnapshot); auto start = std::chrono::high_resolution_clock::now(); @@ -1065,35 +1241,50 @@ void MultibeamSonarSensor::Implementation::ComputeSonarImage() // -------- Sonar calculations -------- // // ------------------------------------------------// - CArray2D P_Beams = NpsGazeboSonar::sonar_calculation_wrapper( - this->pointCloudImage, // cv::Mat& depth_image (the point cloud image) - normal_image, // cv::Mat& normal_image - hPixelSize, // hPixelSize - vPixelSize, // vPixelSize - hFOV, // hFOV - vFOV, // VFOV - hPixelSize, // _beam_azimuthAngleWidth - verticalFOV / 180 * M_PI, // _beam_elevationAngleWidth - hPixelSize, // _ray_azimuthAngleWidth - this->elevation_angles, // _ray_elevationAngles - vPixelSize * (raySkips + 1), // _ray_elevationAngleWidth - this->soundSpeed, // _soundSpeed - this->maxDistance, // _maxDistance - this->sourceLevel, // _sourceLevel - this->nBeams, // _nBeams - this->nRays, // _nRays - this->raySkips, // _raySkips - this->sonarFreq, // _sonarFreq - this->bandwidth, // _bandwidth - this->nFreq, // _nFreq - this->reflectivityImage, // reflectivity_image - this->attenuation, // _attenuation - this->window, // _window - this->beamCorrector, // _beamCorrector - this->beamCorrectorSum, // _beamCorrectorSum - this->debugFlag, // debugFlag - this->blazingFlag // _blazingFlag - ); + SonarComputeInput input; + input.depthImage = &depthSnapshot; + input.normalImage = &normal_image; + input.reflectivityImage = &this->reflectivityImage; + input.nBeams = this->nBeams; + input.nRays = this->nRays; + input.nFreq = this->nFreq; + input.raySkips = this->raySkips; + input.hFOV = this->hFOV; + input.vFOV = this->vFOV; + input.maxDistance = this->maxDistance; + input.soundSpeed = this->soundSpeed; + input.attenuation = this->attenuation; + input.sensorGain = this->sensorGain; + input.debugFlag = this->debugFlag; + input.blazingFlag = this->blazingFlag; + input.window = this->window; + input.rangeVector = this->rangeVector; + input.beamCorrector = this->beamCorrector; + input.beamCorrectorSum = this->beamCorrectorSum; + input.sourceLevel = this->sourceLevel; + input.bandwidth = this->bandwidth; + input.sonarFreq = this->sonarFreq; + input.elevation_angles.assign(this->elevation_angles, this->elevation_angles + this->nRays); + input.frameIndex = this->frameCounter++; + input.seed = this->sonarSeed; + + SonarComputeOutput backendOutput; + if (!this->computeBackend->Compute(input, backendOutput)) + { + RCLCPP_ERROR_STREAM( + this->ros_node_->get_logger(), + "Sonar compute backend failed: " << this->computeBackend->Name()); + return; + } + + CArray2D P_Beams(CArray(this->nFreq), this->nBeams); + for (int beam = 0; beam < this->nBeams; ++beam) + { + for (int freq = 0; freq < this->nFreq; ++freq) + { + P_Beams[beam][freq] = backendOutput.At(beam, freq); + } + } // For calc time measure auto stop = std::chrono::high_resolution_clock::now(); @@ -1101,9 +1292,11 @@ void MultibeamSonarSensor::Implementation::ComputeSonarImage() if (debugFlag) { + double timeMs = duration.count() / 1000.0; RCLCPP_INFO_STREAM( - this->ros_node_->get_logger(), - "GPU Sonar Frame Calc Time " << duration.count() / 10000 << "/100 [s]\n"); + this->ros_node_->get_logger(), this->computeBackend->Name() + << " Sonar Frame Calc Time " << std::fixed + << std::setprecision(3) << timeMs << " [ms]\n"); } // CSV log write stream @@ -1164,15 +1357,26 @@ void MultibeamSonarSensor::Implementation::ComputeSonarImage() } } - rclcpp::Time now = this->ros_node_->now(); - + // Use the single timestamp captured at snapshot time for all publications. + // This ensures sonarRawDataMsg, sonarImgMsg, and normalImgMsg all carry the + // same stamp - the sim time when the depth frame was actually recorded - + // so RViz TF lookups succeed and images don't flicker. + + // FEEDBACK_NEEDED/TODO: Solved? + // Using frameName (from SDF tag) as the ROS frame_id. + // This is user-friendly but requires the SDF value to match an actual TF frame exactly. + // The old approach used frameId (Gazebo's internal sensor ID) which was more likely + // to match the TF tree automatically. + // If TF lookup errors appear in RViz, shall revert back to frameId. this->sonarRawDataMsg = marine_acoustic_msgs::msg::ProjectedSonarImage(); this->sonarRawDataMsg.header.frame_id = this->frameId; - this->sonarRawDataMsg.header.stamp.sec = static_cast(now.seconds()); + this->sonarRawDataMsg.header.frame_id = this->frameName; + + this->sonarRawDataMsg.header.stamp.sec = static_cast(capturedStamp.seconds()); this->sonarRawDataMsg.header.stamp.nanosec = - static_cast(now.nanoseconds() % 1000000000); + static_cast(capturedStamp.nanoseconds() % 1000000000); marine_acoustic_msgs::msg::PingInfo ping_info_msg_; @@ -1209,7 +1413,7 @@ void MultibeamSonarSensor::Implementation::ComputeSonarImage() sonar_image_data.beam_count = this->nBeams; // this->sonar_image_raw_msg_.data_size = 1; // sizeof(float) * nFreq * nBeams; std::vector intensities; - int Intensity[this->nBeams][this->nFreq]; + // int Intensity[this->nBeams][this->nFreq]; for (size_t r = 0; r < P_Beams[0].size(); r++) { @@ -1224,9 +1428,6 @@ void MultibeamSonarSensor::Implementation::ComputeSonarImage() this->sonarRawDataMsg.image = sonar_image_data; this->sonarImageRawPub->publish(this->sonarRawDataMsg); - // Construct visual sonar image for rqt plot in sensor::image msg format - cv_bridge::CvImage img_bridge; - // Generate image of CV_8UC1 cv::Mat Intensity_image = cv::Mat::zeros(cv::Size(this->nBeams, this->nFreq), CV_8UC1); @@ -1310,33 +1511,25 @@ void MultibeamSonarSensor::Implementation::ComputeSonarImage() cv::Mat Itensity_image_color; cv::applyColorMap(Intensity_image, Itensity_image_color, cv::COLORMAP_HOT); - now = this->ros_node_->now(); this->sonarImgMsg.header.frame_id = this->frameName; - this->sonarImgMsg.header.stamp.sec = static_cast(now.seconds()); - this->sonarImgMsg.header.stamp.nanosec = static_cast(now.nanoseconds() % 1000000000); + this->sonarImgMsg.header.stamp.sec = static_cast(capturedStamp.seconds()); + this->sonarImgMsg.header.stamp.nanosec = + static_cast(capturedStamp.nanoseconds() % 1000000000); - img_bridge = cv_bridge::CvImage( - this->sonarImgMsg.header, sensor_msgs::image_encodings::BGR8, Itensity_image_color); - // from cv_bridge to sensor_msgs::Image - img_bridge.toImageMsg(this->sonarImgMsg); + CopyImageToMessage(Itensity_image_color, sensor_msgs::image_encodings::BGR8, this->sonarImgMsg); this->sonarImagePub->publish(this->sonarImgMsg); - // ---------------------------------------- End of sonar calculation - - now = this->ros_node_->now(); this->normalImgMsg.header.frame_id = this->frameName; - this->normalImgMsg.header.stamp.sec = static_cast(now.seconds()); - this->normalImgMsg.header.stamp.nanosec = static_cast(now.nanoseconds() % 1000000000); + this->normalImgMsg.header.stamp.sec = static_cast(capturedStamp.seconds()); + this->normalImgMsg.header.stamp.nanosec = + static_cast(capturedStamp.nanoseconds() % 1000000000); cv::Mat normal_image8; normal_image.convertTo(normal_image8, CV_8UC3, 255.0); - img_bridge = cv_bridge::CvImage( - this->normalImgMsg.header, sensor_msgs::image_encodings::RGB8, normal_image8); - img_bridge.toImageMsg(this->normalImgMsg); - // from cv_bridge to sensor_msgs::Image + CopyImageToMessage(normal_image8, sensor_msgs::image_encodings::RGB8, this->normalImgMsg); this->normalImagePub->publish(this->normalImgMsg); - this->lock_.unlock(); + // ---------------------------------------- End of sonar calculation } } // namespace sensors diff --git a/gazebo/posim_gz_multibeam_sonar/multibeam_sonar/MultibeamSonarSensor.hh b/gazebo/posim_gz_multibeam_sonar/multibeam_sonar/MultibeamSonarSensor.hh index 3138bd10..f8eb8539 100644 --- a/gazebo/posim_gz_multibeam_sonar/multibeam_sonar/MultibeamSonarSensor.hh +++ b/gazebo/posim_gz_multibeam_sonar/multibeam_sonar/MultibeamSonarSensor.hh @@ -3,9 +3,11 @@ #include #include +#include #include #include #include +#include #include #include #include @@ -24,6 +26,7 @@ #include #include "AcousticBeam.hh" #include "AxisAlignedPatch2.hh" +#include "sonar_compute_backend.hh" namespace gz { @@ -108,8 +111,16 @@ private: GZ_UTILS_WARN_IGNORE__DLL_INTERFACE_MISSING mutable std::mutex rayMutex; GZ_UTILS_WARN_RESUME__DLL_INTERFACE_MISSING + + // Background sonar compute thread and synchronization + std::thread computeThread_; + std::condition_variable computeCV_; + std::mutex computeMutex_; + bool newFrameReady_{false}; + bool stopThread_{false}; // ROS node pointer std::shared_ptr ros_node_; + std::shared_ptr ros_executor_; // SDF sensor element sdf::ElementPtr sensorSdf; @@ -169,6 +180,12 @@ private: bool constMu; double mu; + // Backend selection and determinism controls. + std::string requestedBackend{"auto"}; + uint64_t sonarSeed{12345}; + uint64_t frameCounter{0}; + std::unique_ptr computeBackend; + // Beam corrector float beamCorrectorSum; float ** beamCorrector; @@ -200,7 +217,7 @@ private: const std::string & /*_format*/); void ComputeSonarImage(); cv::Mat ComputeNormalImage(cv::Mat & depth); - void ComputeCorrector(); + void ComputeCorrector(int _snapshotWidth, int _nBeams); // Connections gz::common::ConnectionPtr rayConnection; diff --git a/gazebo/posim_gz_multibeam_sonar/multibeam_sonar/hooks/multibeam_sonar.dsv.in b/gazebo/posim_gz_multibeam_sonar/multibeam_sonar/hooks/multibeam_sonar.dsv.in index 07f53dcf..4e8a0b0a 100644 --- a/gazebo/posim_gz_multibeam_sonar/multibeam_sonar/hooks/multibeam_sonar.dsv.in +++ b/gazebo/posim_gz_multibeam_sonar/multibeam_sonar/hooks/multibeam_sonar.dsv.in @@ -1,2 +1,3 @@ prepend-non-duplicate;GZ_SIM_SYSTEM_PLUGIN_PATH;lib/@PROJECT_NAME@/ prepend-non-duplicate;LD_LIBRARY_PATH;lib/@PROJECT_NAME@/ +prepend-non-duplicate;DYLD_LIBRARY_PATH;lib/@PROJECT_NAME@/ diff --git a/gazebo/posim_gz_multibeam_sonar/multibeam_sonar/package.xml b/gazebo/posim_gz_multibeam_sonar/multibeam_sonar/package.xml index 5ebfc72c..6122dd7b 100644 --- a/gazebo/posim_gz_multibeam_sonar/multibeam_sonar/package.xml +++ b/gazebo/posim_gz_multibeam_sonar/multibeam_sonar/package.xml @@ -8,16 +8,16 @@ Gaurav kumar Ye-Seol Kwon Apache 2.0 - ament_cmake + ament_cmake rclcpp sensor_msgs geometry_msgs - cv_bridge image_transport marine_acoustic_msgs pcl_conversions posim_interfaces protobuf + wgpu_vendor ament_lint_auto ament_cmake diff --git a/gazebo/posim_gz_multibeam_sonar/multibeam_sonar/sonar_compute_backend.hh b/gazebo/posim_gz_multibeam_sonar/multibeam_sonar/sonar_compute_backend.hh new file mode 100644 index 00000000..3e2e89ae --- /dev/null +++ b/gazebo/posim_gz_multibeam_sonar/multibeam_sonar/sonar_compute_backend.hh @@ -0,0 +1,84 @@ +#pragma once + +#include +#include +#include +#include +#include + +#include + +namespace gz +{ +namespace sensors +{ + +struct SonarComputeInput +{ + const cv::Mat * depthImage{nullptr}; + const cv::Mat * normalImage{nullptr}; + const cv::Mat * reflectivityImage{nullptr}; + + int nBeams{0}; + int nRays{0}; + int nFreq{0}; + int raySkips{1}; + + double hFOV{0.0}; + double vFOV{0.0}; + double maxDistance{0.0}; + double soundSpeed{1500.0}; + double attenuation{0.0}; + + float sensorGain{0.02f}; + bool debugFlag{false}; + bool blazingFlag{false}; + + const float * window{nullptr}; + const float * rangeVector{nullptr}; + float ** beamCorrector{nullptr}; + float beamCorrectorSum{0.0f}; + + double sourceLevel{220.0}; + double bandwidth{29.5e6}; + double sonarFreq{900e3}; + std::vector elevation_angles; + + uint64_t frameIndex{0}; + uint64_t seed{0}; +}; + +struct SonarComputeOutput +{ + int nBeams{0}; + int nFreq{0}; + std::vector> beamSpectrum; + double computeMicros{0.0}; + + std::complex & At(int beam, int freq) + { + return this->beamSpectrum[static_cast(beam) * this->nFreq + freq]; + } + + const std::complex & At(int beam, int freq) const + { + return this->beamSpectrum[static_cast(beam) * this->nFreq + freq]; + } +}; + +class ComputeBackend +{ +public: + virtual ~ComputeBackend() = default; + + virtual const char * Name() const = 0; + + virtual bool Initialize(const SonarComputeInput & prototype) = 0; + + virtual bool Compute(const SonarComputeInput & input, SonarComputeOutput & output) = 0; +}; + +std::unique_ptr CreateComputeBackend(const std::string & requestedBackend); + +} // namespace sensors +} // namespace gz diff --git a/gazebo/posim_gz_multibeam_sonar/multibeam_sonar/sonar_compute_cpu.cc b/gazebo/posim_gz_multibeam_sonar/multibeam_sonar/sonar_compute_cpu.cc new file mode 100644 index 00000000..b010268b --- /dev/null +++ b/gazebo/posim_gz_multibeam_sonar/multibeam_sonar/sonar_compute_cpu.cc @@ -0,0 +1,234 @@ +#include "sonar_compute_backend.hh" + +#include +#include +#include +#include + +// NOTE: This file serves two responsibilities -- it implements the CPU compute +// backend (CpuComputeBackend) AND hosts the CreateComputeBackend() factory +// that selects between CPU and WGPU at runtime. Because the factory lives here, +// this file must include sonar_compute_wgpu.hh, which creates a hard compile-time +// dependency on the WGPU backend even when only CPU support is needed. +// +// TODO: Move CreateComputeBackend() into its own file (e.g. sonar_compute_factory.cc) +// so that the CPU and WGPU backends are fully decoupled and the factory is the +// only translation unit that needs to know about both. +#include "sonar_compute_cuda.hh" +#include "sonar_compute_wgpu.hh" + +namespace gz +{ +namespace sensors +{ +namespace +{ + +class CpuComputeBackend : public ComputeBackend +{ +public: + const char * Name() const override { return "cpu"; } + + bool Initialize(const SonarComputeInput &) override { return true; } + + bool Compute(const SonarComputeInput & input, SonarComputeOutput & output) override + { + if (!input.depthImage || !input.normalImage || input.nBeams <= 0 || input.nFreq <= 0) + { + return false; + } + + auto start = std::chrono::high_resolution_clock::now(); + + output.nBeams = input.nBeams; + output.nFreq = input.nFreq; + output.beamSpectrum.assign( + static_cast(input.nBeams) * static_cast(input.nFreq), + std::complex(0.0f, 0.0f)); + + const cv::Mat & depth = *input.depthImage; + const cv::Mat & normal = *input.normalImage; + + if (depth.empty() || depth.type() != CV_32FC1 || normal.empty() || normal.type() != CV_32FC3) + { + return false; + } + + const int rows = depth.rows; + const int cols = depth.cols; + const float maxDistance = static_cast(std::max(1e-6, input.maxDistance)); + const float attenuation = static_cast(input.attenuation); + + const float source_term = + std::sqrt(std::pow(10.0f, static_cast(input.sourceLevel) / 10.0f)) * 1e-6f; + const float h_pixel = static_cast(input.hFOV) / std::max(1, input.nBeams - 1); + const float v_pixel = static_cast(input.vFOV) / std::max(1, input.nRays - 1); + const float area_scaler = h_pixel * v_pixel * static_cast(input.raySkips + 1); + + for (int r = 0; r < rows; ++r) + { + const float * depthRow = depth.ptr(r); + const cv::Vec3f * normalRow = normal.ptr(r); + for (int c = 0; c < cols; ++c) + { + const float d = depthRow[c]; + if (!std::isfinite(d) || d <= 0.0f || d > maxDistance) + { + continue; + } + + const int beam = std::min(input.nBeams - 1, (c * input.nBeams) / std::max(1, cols)); + const int bin = + std::min(input.nFreq - 1, static_cast((d / maxDistance) * input.nFreq)); + + const float nz = std::max(0.0f, normalRow[c][2]); + float reflectivity = 1.0f; + if (input.reflectivityImage && !input.reflectivityImage->empty()) + { + reflectivity = input.reflectivityImage->at(r, c); + } + + // CUDA-aligned physics: spread + lambert + stochastic phase term + const float lambert = std::sqrt(nz); // sqrt(cos(incidence)); nz == cos(acos(nz)) + const float spread = 1.0f / std::max(1e-6f, d * d); + const float target_area = std::sqrt(d * area_scaler); + const float amplitude_mag = input.sensorGain * source_term * spread * + std::exp(-2.0f * attenuation * d) * lambert * + std::sqrt(reflectivity) * target_area; + + const float phase_noise = static_cast((beam * input.nRays + r) % input.nFreq) / + static_cast(input.nFreq) * 2.0f * static_cast(M_PI); + output.At(beam, bin) += std::complex( + amplitude_mag * std::cos(phase_noise) / std::sqrt(2.0f), + amplitude_mag * std::sin(phase_noise) / std::sqrt(2.0f)); + } + } + + // Apply the configured range window. + if (input.window) + { + for (int b = 0; b < input.nBeams; ++b) + { + for (int f = 0; f < input.nFreq; ++f) + { + output.At(b, f) *= input.window[f]; + } + } + } + + // Apply beam correction matrix when available. + if (input.beamCorrector && input.beamCorrectorSum > 0.0f) + { + std::vector> corrected = output.beamSpectrum; + const float norm = input.beamCorrectorSum; + for (int b = 0; b < input.nBeams; ++b) + { + for (int f = 0; f < input.nFreq; ++f) + { + std::complex accum{0.0f, 0.0f}; + for (int k = 0; k < input.nBeams; ++k) + { + const float w = input.beamCorrector[b][k] / norm; + accum += output.At(k, f) * w; + } + corrected[static_cast(b) * input.nFreq + f] = accum; + } + } + output.beamSpectrum.swap(corrected); + } + + const auto stop = std::chrono::high_resolution_clock::now(); + output.computeMicros = + std::chrono::duration_cast(stop - start).count(); + return true; + } +}; + +} // namespace + +std::unique_ptr CreateComputeBackend(const std::string & requestedBackend) +{ + std::cerr << "[sonar_compute_factory] CreateComputeBackend called with: " << requestedBackend + << std::endl; + + std::string backend = requestedBackend; + std::transform( + backend.begin(), backend.end(), backend.begin(), + [](unsigned char c) { return static_cast(std::tolower(c)); }); + + // An empty value behaves like an unset variable. + if (backend.empty()) + { + backend = "auto"; + } + + if (backend == "wgpu") + { + std::cerr << "[sonar_compute_factory] Attempting WgpuComputeBackend initialization..." + << std::endl; + auto wgpuBackend = std::make_unique(); + SonarComputeInput probe; + if (wgpuBackend->Initialize(probe)) + { + std::cerr << "[sonar_compute_factory] SUCCESS: WgpuComputeBackend initialized!" << std::endl; + return wgpuBackend; + } + std::cerr << "[sonar_compute_factory] WGPU backend requested but FAILED to initialize." + << std::endl; + return nullptr; + } + + if (backend == "cuda") + { + std::cerr << "[sonar_compute_factory] Attempting CudaComputeBackend initialization..." + << std::endl; + auto cudaBackend = std::make_unique(); + SonarComputeInput probe; + if (cudaBackend->Initialize(probe)) + { + std::cerr << "[sonar_compute_factory] SUCCESS: CudaComputeBackend initialized!" << std::endl; + return cudaBackend; + } + std::cerr << "[sonar_compute_factory] CUDA backend requested but failed to initialize." + << std::endl; + return nullptr; + } + + if (backend == "cpu") + { + std::cerr << "[sonar_compute_factory] Creating CPU backend" << std::endl; + return std::make_unique(); + } + + if (backend == "auto") + { + std::cerr << "[sonar_compute_factory] Auto-selecting best available backend..." << std::endl; + SonarComputeInput probe; + auto cudaBackend = std::make_unique(); + if (cudaBackend->Initialize(probe)) + { + std::cerr << "[sonar_compute_factory] SUCCESS: Auto-selected CUDA backend" << std::endl; + return cudaBackend; + } + + auto wgpuBackend = std::make_unique(); + if (wgpuBackend->Initialize(probe)) + { + std::cerr << "[sonar_compute_factory] SUCCESS: Auto-selected WGPU backend (CUDA unavailable)" + << std::endl; + return wgpuBackend; + } + std::cerr << "[sonar_compute_factory] SUCCESS: Auto-selected CPU backend (both CUDA and WGPU " + "unavailable)" + << std::endl; + return std::make_unique(); + } + + // Reject unknown values rather than silently substituting a different backend. + std::cerr << "[sonar_compute_factory] Unknown sonar compute backend '" << requestedBackend + << "'. Valid values: auto, wgpu, cuda, cpu." << std::endl; + return nullptr; +} + +} // namespace sensors +} // namespace gz diff --git a/gazebo/posim_gz_multibeam_sonar/multibeam_sonar/sonar_compute_cuda.cc b/gazebo/posim_gz_multibeam_sonar/multibeam_sonar/sonar_compute_cuda.cc new file mode 100644 index 00000000..858d6e74 --- /dev/null +++ b/gazebo/posim_gz_multibeam_sonar/multibeam_sonar/sonar_compute_cuda.cc @@ -0,0 +1,100 @@ +#include "sonar_compute_cuda.hh" + +#include +#include + +#ifdef POSIM_HAS_CUDA_BACKEND +#include +#include "sonar_calculation_cuda.cuh" +#endif + +namespace gz +{ +namespace sensors +{ + +CudaComputeBackend::~CudaComputeBackend() +{ +#ifdef POSIM_HAS_CUDA_BACKEND + if (this->initialized) + { + NpsGazeboSonar::free_cuda_memory(); + } +#endif +} + +const char * CudaComputeBackend::Name() const { return "cuda"; } + +bool CudaComputeBackend::Initialize(const SonarComputeInput &) +{ +#ifdef POSIM_HAS_CUDA_BACKEND + int deviceCount = 0; + const cudaError_t err = cudaGetDeviceCount(&deviceCount); + if (err != cudaSuccess || deviceCount <= 0) + { + return false; + } + + this->initialized = true; + return true; +#else + return false; +#endif +} + +bool CudaComputeBackend::Compute(const SonarComputeInput & input, SonarComputeOutput & output) +{ +#ifdef POSIM_HAS_CUDA_BACKEND + if (!this->initialized) + { + return false; + } + + if ( + !input.depthImage || !input.normalImage || !input.reflectivityImage || !input.window || + !input.beamCorrector || input.elevation_angles.empty()) + { + return false; + } + + auto start = std::chrono::high_resolution_clock::now(); + + const double hPixelSize = input.hFOV / static_cast(input.nBeams - 1); + const double vPixelSize = input.vFOV / static_cast(input.nRays - 1); + + const NpsGazeboSonar::CArray2D beams = NpsGazeboSonar::sonar_calculation_wrapper( + *input.depthImage, *input.normalImage, hPixelSize, vPixelSize, input.hFOV, input.vFOV, + hPixelSize, input.vFOV / 180.0 * M_PI, hPixelSize, + const_cast(input.elevation_angles.data()), vPixelSize * (input.raySkips + 1), + input.soundSpeed, input.maxDistance, input.sourceLevel, input.nBeams, input.nRays, + input.raySkips, input.sonarFreq, input.bandwidth, input.nFreq, *input.reflectivityImage, + input.attenuation, const_cast(input.window), input.beamCorrector, + input.beamCorrectorSum, input.debugFlag, input.blazingFlag); + + output.nBeams = input.nBeams; + output.nFreq = input.nFreq; + output.beamSpectrum.assign( + static_cast(input.nBeams) * static_cast(input.nFreq), + std::complex(0.0f, 0.0f)); + + for (int beam = 0; beam < input.nBeams; ++beam) + { + for (int freq = 0; freq < input.nFreq; ++freq) + { + output.At(beam, freq) = beams[beam][freq]; + } + } + + auto stop = std::chrono::high_resolution_clock::now(); + output.computeMicros = + std::chrono::duration_cast(stop - start).count(); + return true; +#else + static_cast(input); + static_cast(output); + return false; +#endif +} + +} // namespace sensors +} // namespace gz diff --git a/gazebo/posim_gz_multibeam_sonar/multibeam_sonar/sonar_compute_cuda.hh b/gazebo/posim_gz_multibeam_sonar/multibeam_sonar/sonar_compute_cuda.hh new file mode 100644 index 00000000..7dfdb47f --- /dev/null +++ b/gazebo/posim_gz_multibeam_sonar/multibeam_sonar/sonar_compute_cuda.hh @@ -0,0 +1,26 @@ +#pragma once + +#include "sonar_compute_backend.hh" + +namespace gz +{ +namespace sensors +{ + +class CudaComputeBackend : public ComputeBackend +{ +public: + ~CudaComputeBackend() override; + + const char * Name() const override; + + bool Initialize(const SonarComputeInput & prototype) override; + + bool Compute(const SonarComputeInput & input, SonarComputeOutput & output) override; + +private: + bool initialized{false}; +}; + +} // namespace sensors +} // namespace gz diff --git a/gazebo/posim_gz_multibeam_sonar/multibeam_sonar/sonar_compute_wgpu.cc b/gazebo/posim_gz_multibeam_sonar/multibeam_sonar/sonar_compute_wgpu.cc new file mode 100644 index 00000000..75bf0f04 --- /dev/null +++ b/gazebo/posim_gz_multibeam_sonar/multibeam_sonar/sonar_compute_wgpu.cc @@ -0,0 +1,179 @@ +#include "sonar_compute_wgpu.hh" + +#include +#include +#include +#include +#include +#include + +extern "C" +{ + float * sonar_wgpu_compute( + const float * depth_flat, const float * normal_flat, const float * reflectivity_flat, + const float * beam_corrector_flat, const float * window_flat, uint32_t n_beams, uint32_t n_rays, + uint32_t n_freq, uint32_t ray_skips, float sound_speed, float max_distance, float source_level, + float attenuation, float sensor_gain, float bandwidth, float beam_corrector_sum, + float area_scaler, float h_fov, float v_fov, uint64_t frame_index, uint64_t seed, + bool debug_flag); + + void sonar_wgpu_free(float * ptr, size_t len); +} + +namespace gz +{ +namespace sensors +{ + +const char * WgpuComputeBackend::Name() const { return "wgpu"; } + +bool WgpuComputeBackend::Initialize(const SonarComputeInput &) +{ + // Probe the Rust FFI to check if a GPU adapter is available. If it is not, report failure + // so the factory can pick another backend instead of silently computing on the CPU. + std::vector depth(1, 0.0f); + std::vector normal(3, 0.0f); + std::vector reflectivity(1, 0.0f); + std::vector beamCorrector(1, 1.0f); + std::vector window(4, 1.0f); + + float * probe = sonar_wgpu_compute( + depth.data(), normal.data(), reflectivity.data(), beamCorrector.data(), window.data(), 1u, 1u, + 4u, 0u, 1500.0f, 10.0f, 220.0f, 0.0f, 1.0f, 29500000.0f, 1.0f, 1e-6f, 0.0f, 0.0f, 0u, 1u, + false); + + if (!probe) + { + std::cerr << "[sonar_wgpu] GPU probe failed: no usable WGPU adapter." << std::endl; + return false; + } + sonar_wgpu_free(probe, 1u * 4u * 2u); + + this->initialized = true; + return true; +} + +bool WgpuComputeBackend::Compute(const SonarComputeInput & input, SonarComputeOutput & output) +{ + if (!this->initialized) + { + return false; + } + + if (!input.depthImage || !input.normalImage || !input.reflectivityImage) + { + return false; + } + + const cv::Mat & depth = *input.depthImage; + const cv::Mat & normal = *input.normalImage; + const cv::Mat & reflectivity = *input.reflectivityImage; + + if (depth.empty() || normal.empty() || reflectivity.empty()) + { + return false; + } + + const int nBeams = input.nBeams; + const int nRays = input.nRays; + const int nFreq = input.nFreq; + + if (nBeams <= 0 || nRays <= 0 || nFreq <= 0) + { + return false; + } + + auto start = std::chrono::high_resolution_clock::now(); + + // Compute area_scaler matching CUDA: ray_azimuthAngleWidth * ray_elevationAngleWidth + // where ray_azimuthAngleWidth = hPixelSize + // ray_elevationAngleWidth = vPixelSize * (raySkips + 1) + const float hPixelSize = static_cast(input.hFOV) / std::max(1, nBeams - 1); + const float vPixelSize = static_cast(input.vFOV) / std::max(1, nRays - 1); + + const int raySkipsFactor = input.raySkips + 1; + const float area_scaler = hPixelSize * vPixelSize * raySkipsFactor; + + // This transposes from OpenCV's [row=ray][col=beam] to [beam][ray] for WGPU. + // Flatten as [beam][ray]. + std::vector depthFlat(static_cast(nBeams) * nRays, 0.0f); + std::vector reflFlat(static_cast(nBeams) * nRays, 0.0f); + std::vector normalFlat(static_cast(nBeams) * nRays * 3, 0.0f); + + for (int beam = 0; beam < nBeams; ++beam) + { + for (int ray = 0; ray < nRays; ++ray) + { + if (ray >= depth.rows || beam >= depth.cols) + { + continue; + } + const size_t idx = static_cast(beam) * nRays + ray; + depthFlat[idx] = depth.at(ray, beam); + reflFlat[idx] = reflectivity.at(ray, beam); + const cv::Vec3f n = normal.at(ray, beam); + normalFlat[idx * 3 + 0] = n[0]; + normalFlat[idx * 3 + 1] = n[1]; + normalFlat[idx * 3 + 2] = n[2]; + } + } + + std::vector beamCorrectorFlat(static_cast(nBeams) * nBeams, 0.0f); + if (input.beamCorrector) + { + for (int r = 0; r < nBeams; ++r) + { + for (int c = 0; c < nBeams; ++c) + { + beamCorrectorFlat[static_cast(r) * nBeams + c] = input.beamCorrector[r][c]; + } + } + } + + std::vector windowFlat(static_cast(nFreq), 1.0f); + if (input.window) + { + for (int i = 0; i < nFreq; ++i) + { + windowFlat[static_cast(i)] = input.window[i]; + } + } + + float * result = sonar_wgpu_compute( + depthFlat.data(), normalFlat.data(), reflFlat.data(), beamCorrectorFlat.data(), + windowFlat.data(), static_cast(nBeams), static_cast(nRays), + static_cast(nFreq), static_cast(std::max(0, input.raySkips)), + static_cast(input.soundSpeed), static_cast(input.maxDistance), + static_cast(input.sourceLevel), static_cast(input.attenuation), input.sensorGain, + static_cast(input.bandwidth), input.beamCorrectorSum, area_scaler, + static_cast(input.hFOV), static_cast(input.vFOV), input.frameIndex, input.seed, + input.debugFlag); + + if (!result) + { + return false; + } + + output.nBeams = nBeams; + output.nFreq = nFreq; + output.beamSpectrum.assign(static_cast(nBeams) * nFreq, std::complex(0.0f, 0.0f)); + + for (int b = 0; b < nBeams; ++b) + { + for (int f = 0; f < nFreq; ++f) + { + const size_t base = (static_cast(b) * nFreq + f) * 2; + output.At(b, f) = std::complex(result[base], result[base + 1]); + } + } + + sonar_wgpu_free(result, static_cast(nBeams) * nFreq * 2); + + auto stop = std::chrono::high_resolution_clock::now(); + output.computeMicros = static_cast( + std::chrono::duration_cast(stop - start).count()); + return true; +} + +} // namespace sensors +} // namespace gz diff --git a/gazebo/posim_gz_multibeam_sonar/multibeam_sonar/sonar_compute_wgpu.hh b/gazebo/posim_gz_multibeam_sonar/multibeam_sonar/sonar_compute_wgpu.hh new file mode 100644 index 00000000..8929a6db --- /dev/null +++ b/gazebo/posim_gz_multibeam_sonar/multibeam_sonar/sonar_compute_wgpu.hh @@ -0,0 +1,26 @@ +#pragma once + +#include + +#include "sonar_compute_backend.hh" + +namespace gz +{ +namespace sensors +{ + +class WgpuComputeBackend : public ComputeBackend +{ +public: + const char * Name() const override; + + bool Initialize(const SonarComputeInput & prototype) override; + + bool Compute(const SonarComputeInput & input, SonarComputeOutput & output) override; + +private: + bool initialized{false}; +}; + +} // namespace sensors +} // namespace gz diff --git a/gazebo/posim_gz_multibeam_sonar/multibeam_sonar_demo/CMakeLists.txt b/gazebo/posim_gz_multibeam_sonar/multibeam_sonar_demo/CMakeLists.txt index 1bf592ad..6c8704dd 100644 --- a/gazebo/posim_gz_multibeam_sonar/multibeam_sonar_demo/CMakeLists.txt +++ b/gazebo/posim_gz_multibeam_sonar/multibeam_sonar_demo/CMakeLists.txt @@ -2,24 +2,24 @@ cmake_minimum_required(VERSION 3.8) project(posim_multibeam_sonar_demo) find_package(ament_cmake REQUIRED) -find_package(CUDAToolkit QUIET) -if(CUDAToolkit_FOUND) - install( - DIRECTORY launch - DESTINATION share/${PROJECT_NAME} - ) +install( + DIRECTORY launch + DESTINATION share/${PROJECT_NAME} +) - install( - DIRECTORY rviz - DESTINATION share/${PROJECT_NAME} - ) +install( + DIRECTORY rviz + DESTINATION share/${PROJECT_NAME} +) - ament_environment_hooks( - "${CMAKE_CURRENT_SOURCE_DIR}/hooks/${PROJECT_NAME}.dsv.in") +install( + PROGRAMS + scripts/plotdata.py + DESTINATION lib/${PROJECT_NAME} +) -else() - message(STATUS "CUDA Toolkit not found: Skipping CUDA-specific targets") -endif() +ament_environment_hooks( + "${CMAKE_CURRENT_SOURCE_DIR}/hooks/${PROJECT_NAME}.dsv.in") ament_package() diff --git a/gazebo/posim_gz_multibeam_sonar/multibeam_sonar_demo/launch/multibeam_sonar_demo.launch.py b/gazebo/posim_gz_multibeam_sonar/multibeam_sonar_demo/launch/multibeam_sonar_demo.launch.py index 2554c936..03f49df2 100644 --- a/gazebo/posim_gz_multibeam_sonar/multibeam_sonar_demo/launch/multibeam_sonar_demo.launch.py +++ b/gazebo/posim_gz_multibeam_sonar/multibeam_sonar_demo/launch/multibeam_sonar_demo.launch.py @@ -19,6 +19,7 @@ from launch import LaunchDescription from launch.actions import DeclareLaunchArgument from launch.actions import IncludeLaunchDescription +from launch.actions import SetEnvironmentVariable from launch.conditions import IfCondition from launch.launch_description_sources import PythonLaunchDescriptionSource from launch.substitutions import LaunchConfiguration @@ -27,6 +28,7 @@ def generate_launch_description(): + compute_backend = LaunchConfiguration("compute_backend") pkg_posim_demos = get_package_share_directory("posim_demos") multibeam_sonar_sim = IncludeLaunchDescription( PythonLaunchDescriptionSource( @@ -39,6 +41,7 @@ def generate_launch_description(): "x": "5.8", "z": "2", "yaw": "3.14", + "compute_backend": compute_backend, }.items(), ) @@ -56,8 +59,14 @@ def generate_launch_description(): return LaunchDescription( [ - multibeam_sonar_sim, + DeclareLaunchArgument( + "compute_backend", + default_value="auto", + description="Sonar backend selection: auto|wgpu|cuda|cpu", + ), DeclareLaunchArgument("rviz", default_value="true", description="Open RViz."), + SetEnvironmentVariable("POSIM_SONAR_COMPUTE_BACKEND", compute_backend), + multibeam_sonar_sim, rviz, ] ) diff --git a/gazebo/posim_gz_multibeam_sonar/multibeam_sonar_demo/package.xml b/gazebo/posim_gz_multibeam_sonar/multibeam_sonar_demo/package.xml index f9ec437d..c9f74208 100644 --- a/gazebo/posim_gz_multibeam_sonar/multibeam_sonar_demo/package.xml +++ b/gazebo/posim_gz_multibeam_sonar/multibeam_sonar_demo/package.xml @@ -7,6 +7,8 @@ Ye-Seol Kwon Apache 2.0 ament_cmake + python3-numpy + python3-matplotlib ament_cmake diff --git a/gazebo/posim_gz_multibeam_sonar/multibeam_sonar_demo/scripts/plotdata.py b/gazebo/posim_gz_multibeam_sonar/multibeam_sonar_demo/scripts/plotdata.py index 3929e443..f3ad4248 100644 --- a/gazebo/posim_gz_multibeam_sonar/multibeam_sonar_demo/scripts/plotdata.py +++ b/gazebo/posim_gz_multibeam_sonar/multibeam_sonar_demo/scripts/plotdata.py @@ -1,3 +1,5 @@ +import os + import numpy as np import matplotlib.pyplot as plt @@ -10,11 +12,17 @@ maxRange = 5 xPlotRange = 10 yPlotRange = xPlotRange * np.cos(45 * np.pi / 180) -filename = "SonarRawData_000001.csv" -beam_angle_file = "SonarRawData_beam_angles.csv" + +results_dir = os.path.join(os.path.dirname(__file__), "..", "..", "..", "..", "results") +os.makedirs(results_dir, exist_ok=True) + +filename = os.path.join(results_dir, "SonarRawData_000001.csv") +beam_angle_file = os.path.join(results_dir, "SonarRawData_beam_angles.csv") bw = 29.9e3 plotSkips = 1 epsilon = 1e-10 +scatter_output_file = os.path.join(results_dir, "SonarScatter.png") +beam_output_file = os.path.join(results_dir, "SonarBeamProfiles.png") # ------------------------------------------- @@ -71,7 +79,8 @@ def load_complex_csv(filepath): plt.ylim(1.02 * np.array([-yPlotRange, yPlotRange])) plt.gca().set_facecolor("k") plt.tight_layout() -plt.show() +plt.savefig(scatter_output_file, dpi=200) +plt.close() # ------------------------------------------- # LINE PLOTS FOR INDIVIDUAL BEAMS @@ -91,4 +100,8 @@ def load_complex_csv(filepath): plt.suptitle("Beam Profiles") plt.tight_layout() -plt.show() +plt.savefig(beam_output_file, dpi=200) +plt.close() + +print(f"Saved {scatter_output_file}") +print(f"Saved {beam_output_file}") diff --git a/gazebo/posim_gz_multibeam_sonar/multibeam_sonar_system/CMakeLists.txt b/gazebo/posim_gz_multibeam_sonar/multibeam_sonar_system/CMakeLists.txt index a601f6a1..5c6b73e6 100644 --- a/gazebo/posim_gz_multibeam_sonar/multibeam_sonar_system/CMakeLists.txt +++ b/gazebo/posim_gz_multibeam_sonar/multibeam_sonar_system/CMakeLists.txt @@ -2,53 +2,42 @@ cmake_minimum_required(VERSION 3.11.0 FATAL_ERROR) project(multibeam_sonar_system) find_package(ament_cmake REQUIRED) -find_package(CUDAToolkit QUIET) - -if(CUDAToolkit_FOUND) - find_package(tinyxml2 CONFIG REQUIRED) - find_package(gz-cmake 5 REQUIRED) - find_package(gz-common 7 REQUIRED COMPONENTS profiler) - find_package(gz-plugin 4 REQUIRED COMPONENTS register) - find_package(gz-rendering 10 REQUIRED OPTIONAL_COMPONENTS ogre ogre2) - find_package(gz-sim 10 REQUIRED) - find_package(gz-sensors 10 REQUIRED) - - if(TARGET gz-rendering::ogre) - set(HAVE_OGRE TRUE) - add_definitions(-DWITH_OGRE) - endif() - - if(TARGET gz-rendering::ogre2) - set(HAVE_OGRE2 TRUE) - add_definitions(-DWITH_OGRE2) - endif() - - add_subdirectory(../multibeam_sonar ${CMAKE_BINARY_DIR}/multibeam_sonar) - - add_library(${PROJECT_NAME} SHARED MultibeamSonarSystem.cc) - target_link_libraries(${PROJECT_NAME} - gz-common::gz-common - gz-plugin::gz-plugin - gz-sim::gz-sim - gz-sensors::gz-sensors - multibeam_sonar - ) - target_include_directories(${PROJECT_NAME} - PUBLIC - ../multibeam_sonar - ) - # Install targets - install(TARGETS ${PROJECT_NAME} - DESTINATION lib/${PROJECT_NAME} - ) - - # Environment hooks - ament_environment_hooks( - "${CMAKE_CURRENT_SOURCE_DIR}/hooks/${PROJECT_NAME}.dsv.in" - ) - -else() - message(STATUS "CUDA Toolkit not found: Skipping CUDA-specific targets") +find_package(gz-cmake 5 REQUIRED) +find_package(gz-common 7 REQUIRED COMPONENTS profiler) +find_package(gz-plugin 4 REQUIRED COMPONENTS register) +find_package(gz-rendering 10 REQUIRED OPTIONAL_COMPONENTS ogre ogre2) +find_package(gz-sim 10 REQUIRED) +find_package(gz-sensors 10 REQUIRED) + +if(TARGET gz-rendering::ogre) + add_definitions(-DWITH_OGRE) endif() +if(TARGET gz-rendering::ogre2) + add_definitions(-DWITH_OGRE2) +endif() + +add_subdirectory(../multibeam_sonar ${CMAKE_BINARY_DIR}/multibeam_sonar) + +add_library(${PROJECT_NAME} SHARED MultibeamSonarSystem.cc) +target_link_libraries(${PROJECT_NAME} + gz-common::gz-common + gz-plugin::gz-plugin + gz-sim::gz-sim + gz-sensors::gz-sensors + multibeam_sonar +) +target_include_directories(${PROJECT_NAME} + PUBLIC + ../multibeam_sonar +) + +install(TARGETS ${PROJECT_NAME} + DESTINATION lib/${PROJECT_NAME} +) + +ament_environment_hooks( + "${CMAKE_CURRENT_SOURCE_DIR}/hooks/${PROJECT_NAME}.dsv.in" +) + ament_package() diff --git a/gazebo/posim_gz_multibeam_sonar/multibeam_sonar_system/MultibeamSonarSystem.cc b/gazebo/posim_gz_multibeam_sonar/multibeam_sonar_system/MultibeamSonarSystem.cc index 81c91798..dc4dd846 100644 --- a/gazebo/posim_gz_multibeam_sonar/multibeam_sonar_system/MultibeamSonarSystem.cc +++ b/gazebo/posim_gz_multibeam_sonar/multibeam_sonar_system/MultibeamSonarSystem.cc @@ -467,6 +467,19 @@ void MultibeamSonarSystem::Implementation::Handle(requests::CreateSensor _reques sensor->SetScene(this->scene); sensor->SetManualSceneUpdate(true); + // SetScene() initializes the sensor; on failure it has no rendering sensor to attach. + if (sensor->RenderingSensors().empty()) + { + gzerr << "Sensor [" << _request.sdf.Name() << "] failed to initialize and was " + << "removed. Other sensors and the simulation keep running." << std::endl; + if (!this->sensorManager.Remove(sensor->Id())) + { + gzerr << "Internal error, missing sensor " + << "[" << _request.sdf.Name() << "]" << std::endl; + } + return; + } + if (this->latestWorldState) { sensor->SetWorldState(*this->latestWorldState); diff --git a/gazebo/posim_gz_multibeam_sonar/multibeam_sonar_system/hooks/multibeam_sonar_system.dsv.in b/gazebo/posim_gz_multibeam_sonar/multibeam_sonar_system/hooks/multibeam_sonar_system.dsv.in index 07f53dcf..1af9a9a4 100644 --- a/gazebo/posim_gz_multibeam_sonar/multibeam_sonar_system/hooks/multibeam_sonar_system.dsv.in +++ b/gazebo/posim_gz_multibeam_sonar/multibeam_sonar_system/hooks/multibeam_sonar_system.dsv.in @@ -1,2 +1,5 @@ prepend-non-duplicate;GZ_SIM_SYSTEM_PLUGIN_PATH;lib/@PROJECT_NAME@/ prepend-non-duplicate;LD_LIBRARY_PATH;lib/@PROJECT_NAME@/ +prepend-non-duplicate;LD_LIBRARY_PATH;lib/multibeam_sonar/ +prepend-non-duplicate;DYLD_LIBRARY_PATH;lib/@PROJECT_NAME@/ +prepend-non-duplicate;DYLD_LIBRARY_PATH;lib/multibeam_sonar/ diff --git a/gazebo/posim_gz_multibeam_sonar/multibeam_sonar_system/package.xml b/gazebo/posim_gz_multibeam_sonar/multibeam_sonar_system/package.xml index 36ba878a..88c68fa2 100644 --- a/gazebo/posim_gz_multibeam_sonar/multibeam_sonar_system/package.xml +++ b/gazebo/posim_gz_multibeam_sonar/multibeam_sonar_system/package.xml @@ -9,7 +9,8 @@ Ye-Seol Kwon Apache 2.0 geometry_msgs - ament_cmake + ament_cmake + multibeam_sonar ament_lint_auto ament_cmake diff --git a/gazebo/posim_gz_multibeam_sonar/validation/wgpu_range_axis/README.md b/gazebo/posim_gz_multibeam_sonar/validation/wgpu_range_axis/README.md new file mode 100644 index 00000000..43e79fe0 --- /dev/null +++ b/gazebo/posim_gz_multibeam_sonar/validation/wgpu_range_axis/README.md @@ -0,0 +1,94 @@ +# WGPU padded-FFT range-axis validation + +This note records the evidence for the range-preserving writeback candidate in +PR #44. The measurements were taken from the PR head at +`6aef91c823af5da073329b84ba617b572965e79e` and from the candidate shader change +applied on top of that commit. + +## Problem + +The WGPU path pads `n_freq` to the next power of two before running its FFT. For +the default sonar configuration, `n_freq = 399` and `padded_n = 512`. The +original shader copied `smem[i]` directly into output bin `i`, while the ROS +message continued to use the original 399-bin physical range vector. + +The expected error signature is therefore + +```text +reported range = physical range * padded_n / n_freq +``` + +For 399 -> 512, the scale is `1.283208`. With a 10 m configured maximum range, +the expected far-range clipping threshold is `10 * 399 / 512 = 7.792969 m`. + +## Candidate + +The candidate keeps the zero-padded power-of-two FFT. During writeback, output +bin `i` samples the complex padded spectrum at `i * padded_n / n_freq` and uses +linear interpolation between the adjacent samples. The output buffer and +published physical range vector remain `n_freq` bins long. + +## Controlled Gazebo experiment + +- Platform: Apple M2, Metal 3, macOS 15.7.3 +- Software: ROS 2 Lyrical, Gazebo Sim 10.5.0 +- Target: one planar panel with its front face placed at the commanded range +- Range: 2, 4, 6, and 8 m +- Incidence angle: 0, 15, 30, and 45 degrees +- Samples: five frames for each backend and condition +- Oracle: an independent PointCloud range measurement checked target placement +- Boundary exclusion: 10 m was excluded because the PointCloud oracle was not + available at the configured maximum range + +The peak range was extracted from the median profile over the five collected +frames. RMSE, MAE, bias, and maximum absolute error were then calculated against +the commanded front-face range. + +| Backend | RMSE (m) | MAE (m) | Max abs. error (m) | +| --- | ---: | ---: | ---: | +| PR #44 CPU | 0.0134 | 0.0111 | 0.0268 | +| Original PR #44 WGPU | 1.4532 | 1.3587 | 1.9582 | +| Candidate WGPU writeback | 0.0540 | 0.0449 | 0.1271 | + +The candidate reduced range RMSE by `96.28%` relative to the original WGPU +writeback in this test matrix. + +![Range-axis validation](rangefix_comparison.png) + +The 16 paired conditions and their measured values are available in +[`rangefix_paired_accuracy.csv`](rangefix_paired_accuracy.csv). + +## Range-bin checks + +The same mechanism was checked using other non-power-of-two range-bin counts. +These are four-distance exploratory checks rather than the full angle matrix. + +| `n_freq` -> FFT size | Original RMSE (m) | Candidate RMSE (m) | +| --- | ---: | ---: | +| 301 -> 512 | 4.7419 | 0.0616 | +| 399 -> 512 | 1.4645 | 0.0385 | +| 511 -> 512 | 0.0397 | 0.0562 | +| 513 -> 1024 | 4.9812 | 0.0110 | + +An unmodified `512 -> 512` power-of-two control produced `0.0252 m` RMSE, +supporting padding-grid mismatch as the dominant mechanism in these controlled +tests. + +![Range-bin RMSE](rangebin_rmse.png) + +The range-bin summaries are available in +[`rangebin_summary.csv`](rangebin_summary.csv). + +## Current limitations + +- NVIDIA CUDA equivalence has not yet been tested. +- The `511 -> 512` case did not benefit from linear interpolation. A no-op or + adaptive policy near a unity padding ratio should be reviewed before treating + this candidate as a universal solution. +- The evidence is limited to controlled Gazebo targets on Apple Metal. It does + not establish HIL, physical-sonar, or real-ocean accuracy. +- End-to-end timing samples overlap; they do not isolate the FFT writeback cost. + +The candidate should therefore be reviewed as a focused correction for the +observed padded-grid range mapping, not as a claim of complete backend +equivalence. diff --git a/gazebo/posim_gz_multibeam_sonar/validation/wgpu_range_axis/rangebin_rmse.png b/gazebo/posim_gz_multibeam_sonar/validation/wgpu_range_axis/rangebin_rmse.png new file mode 100644 index 00000000..5ea98969 Binary files /dev/null and b/gazebo/posim_gz_multibeam_sonar/validation/wgpu_range_axis/rangebin_rmse.png differ diff --git a/gazebo/posim_gz_multibeam_sonar/validation/wgpu_range_axis/rangebin_summary.csv b/gazebo/posim_gz_multibeam_sonar/validation/wgpu_range_axis/rangebin_summary.csv new file mode 100644 index 00000000..ff6e7004 --- /dev/null +++ b/gazebo/posim_gz_multibeam_sonar/validation/wgpu_range_axis/rangebin_summary.csv @@ -0,0 +1,9 @@ +n_freq,padded_n,method,conditions,rmse_m,mae_m,bias_m,max_abs_error_m,unfixed_observable_limit_m,model_valid_conditions,model_residual_rmse_m +301,512,original,4,4.741873322498592,4.072062313556671,0.07206231355667114,8.0,5.87890625,2,0.09010430538758143 +301,512,linear,4,0.06162741057324949,0.04102003574371338,0.03880274295806885,0.11973428726196289,5.87890625,, +399,512,original,4,1.4645263599761196,1.3712374567985535,1.3712374567985535,1.9581937789916992,7.79296875,3,0.054589027573644026 +399,512,linear,4,0.038452310449590285,0.03553497791290283,0.03553497791290283,0.05183887481689453,7.79296875,, +511,512,original,4,0.03973169547237621,0.029047220945358276,0.027741700410842896,0.07310724258422852,9.98046875,4,0.0334465373876441 +511,512,linear,4,0.056245733426789915,0.04862925410270691,0.04732373356819153,0.07310724258422852,9.98046875,, +513,1024,original,4,4.9811570233532825,4.4811443402431905,0.49089726665988564,7.98049414716661,5.009765625,2,0.006504525225154364 +513,1024,linear,4,0.01097666912949125,0.009427905082702637,0.008127570152282715,0.018205642700195312,5.009765625,, diff --git a/gazebo/posim_gz_multibeam_sonar/validation/wgpu_range_axis/rangefix_comparison.png b/gazebo/posim_gz_multibeam_sonar/validation/wgpu_range_axis/rangefix_comparison.png new file mode 100644 index 00000000..3f28a2c0 Binary files /dev/null and b/gazebo/posim_gz_multibeam_sonar/validation/wgpu_range_axis/rangefix_comparison.png differ diff --git a/gazebo/posim_gz_multibeam_sonar/validation/wgpu_range_axis/rangefix_paired_accuracy.csv b/gazebo/posim_gz_multibeam_sonar/validation/wgpu_range_axis/rangefix_paired_accuracy.csv new file mode 100644 index 00000000..20edcd3b --- /dev/null +++ b/gazebo/posim_gz_multibeam_sonar/validation/wgpu_range_axis/rangefix_paired_accuracy.csv @@ -0,0 +1,17 @@ +ground_truth_range_m,incidence_angle_deg,pointcloud_oracle_m,cpu_peak_range_m,wgpu_original_peak_range_m,wgpu_fixed_peak_range_m,cpu_error_m,wgpu_original_error_m,wgpu_fixed_error_m,original_to_fixed_abs_error_reduction_m +2.0,0.0,2.000008106232331,1.9816054105758667,2.6086957454681396,2.0317726135253906,-0.0183945894241333,0.6086957454681396,0.031772613525390625,0.576923131942749 +2.0,15.0,2.001055717468904,1.9816054105758667,2.5836119651794434,2.0317726135253906,-0.0183945894241333,0.5836119651794434,0.031772613525390625,0.5518393516540527 +2.0,30.0,2.00226426124637,1.9816054105758667,2.5836119651794434,2.0066888332366943,-0.0183945894241333,0.5836119651794434,0.006688833236694336,0.576923131942749 +2.0,45.0,2.0039196014410727,2.0066888332366943,2.5836119651794434,2.0317726135253906,0.006688833236694336,0.5836119651794434,0.031772613525390625,0.5518393516540527 +4.0,0.0,4.000015735627504,3.9882943630218506,5.217391490936279,4.013377666473389,-0.011705636978149414,1.2173914909362793,0.013377666473388672,1.2040138244628906 +4.0,15.0,4.002110481263491,3.9882943630218506,5.217391490936279,4.088628768920898,-0.011705636978149414,1.2173914909362793,0.08862876892089844,1.1287627220153809 +4.0,30.0,4.004532337190006,3.9882943630218506,5.192307472229004,4.063545227050781,-0.011705636978149414,1.192307472229004,0.06354522705078125,1.1287622451782227 +4.0,45.0,4.007824897767399,4.013377666473389,5.192307472229004,4.063545227050781,0.013377666473388672,1.192307472229004,0.06354522705078125,1.1287622451782227 +6.0,0.0,6.000023841859836,5.994983196258545,7.700668811798096,6.0451507568359375,-0.005016803741455078,1.7006688117980957,0.0451507568359375,1.6555180549621582 +6.0,15.0,6.003164291383763,5.994983196258545,7.75083589553833,6.020066738128662,-0.005016803741455078,1.75083589553833,0.02006673812866211,1.730769157409668 +6.0,30.0,6.006785392763159,5.994983196258545,7.725752353668213,6.020066738128662,-0.005016803741455078,1.725752353668213,0.02006673812866211,1.7056856155395508 +6.0,45.0,6.011794567110083,6.020066738128662,7.725752353668213,6.020066738128662,0.02006673812866211,1.725752353668213,0.02006673812866211,1.7056856155395508 +8.0,0.0,8.000031471255008,8.00167179107666,9.9581937789917,8.051838874816895,0.0016717910766601562,1.9581937789916992,0.05183887481689453,1.9063549041748047 +8.0,15.0,8.004215240481084,8.00167179107666,9.933110237121582,8.051838874816895,0.0016717910766601562,1.933110237121582,0.05183887481689453,1.8812713623046875 +8.0,30.0,8.009082794192024,8.00167179107666,9.9581937789917,8.051838874816895,0.0016717910766601562,1.9581937789916992,0.05183887481689453,1.9063549041748047 +8.0,45.0,8.015708923342416,8.026755332946777,9.807692527770996,8.127090454101562,0.026755332946777344,1.807692527770996,0.1270904541015625,1.6806020736694336 diff --git a/gazebo/posim_gz_multibeam_sonar/wgpu_vendor/CMakeLists.txt b/gazebo/posim_gz_multibeam_sonar/wgpu_vendor/CMakeLists.txt new file mode 100644 index 00000000..a7b98525 --- /dev/null +++ b/gazebo/posim_gz_multibeam_sonar/wgpu_vendor/CMakeLists.txt @@ -0,0 +1,62 @@ +cmake_minimum_required(VERSION 3.11) +project(wgpu_vendor) + +find_package(ament_cmake REQUIRED) + +# WGPU can use Metal, DirectX, or Vulkan +find_package(Vulkan QUIET) + +find_program(CARGO_EXECUTABLE cargo REQUIRED) + +set(RUST_CRATE_DIR ${CMAKE_CURRENT_SOURCE_DIR}/sonar_wgpu_rust) +set(RUST_TARGET_DIR ${CMAKE_CURRENT_BINARY_DIR}/cargo_target) +set(RUST_PROFILE_DIR ${RUST_TARGET_DIR}/release) +set(RUST_LIB ${RUST_PROFILE_DIR}/libsonar_wgpu.a) + +# Track all Rust source files so colcon/CMake triggers cargo when they change. +file(GLOB_RECURSE RUST_SOURCES + ${RUST_CRATE_DIR}/src/*.rs + ${RUST_CRATE_DIR}/src/shaders/*.wgsl + ${RUST_CRATE_DIR}/Cargo.toml + ${RUST_CRATE_DIR}/Cargo.lock +) + +add_custom_command( + OUTPUT ${RUST_LIB} + COMMAND ${CARGO_EXECUTABLE} build --release --target-dir ${RUST_TARGET_DIR} + WORKING_DIRECTORY ${RUST_CRATE_DIR} + DEPENDS ${RUST_SOURCES} + COMMENT "Building sonar_wgpu Rust static library" + VERBATIM +) + +add_custom_target(sonar_wgpu_rust ALL + DEPENDS ${RUST_LIB} + COMMENT "Ensuring sonar_wgpu Rust static library is built" +) + +add_library(wgpu_vendor INTERFACE) +add_dependencies(wgpu_vendor sonar_wgpu_rust) + +# Export the installed archive, not an imported target from this build tree. +target_link_libraries(wgpu_vendor INTERFACE + "$" + "$/lib/libsonar_wgpu.a>" + dl pthread m +) +if(Vulkan_FOUND) + target_link_libraries(wgpu_vendor INTERFACE Vulkan::Vulkan) +endif() + +add_library(wgpu::wgpu_vendor ALIAS wgpu_vendor) + +install(FILES ${RUST_LIB} DESTINATION lib) + +install(TARGETS wgpu_vendor EXPORT wgpu_vendorTargets) + +# Match the public wgpu::wgpu_vendor target used by the sonar sensor. +ament_export_targets(wgpu_vendorTargets HAS_LIBRARY_TARGET NAMESPACE wgpu::) +if(Vulkan_FOUND) + ament_export_dependencies(Vulkan) +endif() +ament_package() diff --git a/gazebo/posim_gz_multibeam_sonar/wgpu_vendor/package.xml b/gazebo/posim_gz_multibeam_sonar/wgpu_vendor/package.xml new file mode 100644 index 00000000..fd833a36 --- /dev/null +++ b/gazebo/posim_gz_multibeam_sonar/wgpu_vendor/package.xml @@ -0,0 +1,14 @@ + + + wgpu_vendor + 0.1.0 + Vendor package exposing wgpu CMake target for the POSIM sonar backend. + DAVE Maintainers + Apache-2.0 + ament_cmake + cargo + rustc + + ament_cmake + + \ No newline at end of file diff --git a/gazebo/posim_gz_multibeam_sonar/wgpu_vendor/sonar_wgpu_rust/Cargo.lock b/gazebo/posim_gz_multibeam_sonar/wgpu_vendor/sonar_wgpu_rust/Cargo.lock new file mode 100644 index 00000000..3eb292fe --- /dev/null +++ b/gazebo/posim_gz_multibeam_sonar/wgpu_vendor/sonar_wgpu_rust/Cargo.lock @@ -0,0 +1,1104 @@ +# This file is automatically @generated by Cargo. +# It is not intended for manual editing. +version = 4 + +[[package]] +name = "android_system_properties" +version = "0.1.5" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "819e7219dbd41043ac279b19830f2efc897156490d7fd6ea916720117ee66311" +dependencies = [ + "libc", +] + +[[package]] +name = "arrayvec" +version = "0.7.6" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "7c02d123df017efcdfbd739ef81735b36c5ba83ec3c59c80a9d7ecc718f92e50" + +[[package]] +name = "ash" +version = "0.37.3+1.3.251" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "39e9c3835d686b0a6084ab4234fcd1b07dbf6e4767dce60874b12356a25ecd4a" +dependencies = [ + "libloading 0.7.4", +] + +[[package]] +name = "autocfg" +version = "1.5.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "c08606f8c3cbf4ce6ec8e28fb0014a2c086708fe954eaa885384a6165172e7e8" + +[[package]] +name = "bit-set" +version = "0.5.3" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "0700ddab506f33b20a03b13996eccd309a48e5ff77d0d95926aa0210fb4e95f1" +dependencies = [ + "bit-vec", +] + +[[package]] +name = "bit-vec" +version = "0.6.3" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "349f9b6a179ed607305526ca489b34ad0a41aed5f7980fa90eb03160b69598fb" + +[[package]] +name = "bitflags" +version = "1.3.2" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "bef38d45163c2f1dde094a7dfd33ccf595c92905c8f8f4fdc18d06fb1037718a" + +[[package]] +name = "bitflags" +version = "2.11.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "843867be96c8daad0d758b57df9392b6d8d271134fce549de6ce169ff98a92af" + +[[package]] +name = "block" +version = "0.1.6" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "0d8c1fef690941d3e7788d328517591fecc684c084084702d6ff1641e993699a" + +[[package]] +name = "bumpalo" +version = "3.20.2" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "5d20789868f4b01b2f2caec9f5c4e0213b41e3e5702a50157d699ae31ced2fcb" + +[[package]] +name = "bytemuck" +version = "1.25.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "c8efb64bd706a16a1bdde310ae86b351e4d21550d98d056f22f8a7f7a2183fec" +dependencies = [ + "bytemuck_derive", +] + +[[package]] +name = "bytemuck_derive" +version = "1.10.2" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "f9abbd1bc6865053c427f7198e6af43bfdedc55ab791faed4fbd361d789575ff" +dependencies = [ + "proc-macro2", + "quote", + "syn 2.0.117", +] + +[[package]] +name = "cfg-if" +version = "1.0.4" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "9330f8b2ff13f34540b44e946ef35111825727b38d33286ef986142615121801" + +[[package]] +name = "cfg_aliases" +version = "0.1.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "fd16c4719339c4530435d38e511904438d07cce7950afa3718a84ac36c10e89e" + +[[package]] +name = "codespan-reporting" +version = "0.11.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "3538270d33cc669650c4b093848450d380def10c331d38c768e34cac80576e6e" +dependencies = [ + "termcolor", + "unicode-width", +] + +[[package]] +name = "com" +version = "0.6.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "7e17887fd17353b65b1b2ef1c526c83e26cd72e74f598a8dc1bee13a48f3d9f6" +dependencies = [ + "com_macros", +] + +[[package]] +name = "com_macros" +version = "0.6.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "d375883580a668c7481ea6631fc1a8863e33cc335bf56bfad8d7e6d4b04b13a5" +dependencies = [ + "com_macros_support", + "proc-macro2", + "syn 1.0.109", +] + +[[package]] +name = "com_macros_support" +version = "0.6.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "ad899a1087a9296d5644792d7cb72b8e34c1bec8e7d4fbc002230169a6e8710c" +dependencies = [ + "proc-macro2", + "quote", + "syn 1.0.109", +] + +[[package]] +name = "core-foundation" +version = "0.9.4" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "91e195e091a93c46f7102ec7818a2aa394e1e1771c3ab4825963fa03e45afb8f" +dependencies = [ + "core-foundation-sys", + "libc", +] + +[[package]] +name = "core-foundation-sys" +version = "0.8.7" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "773648b94d0e5d620f64f280777445740e61fe701025087ec8b57f45c791888b" + +[[package]] +name = "core-graphics-types" +version = "0.1.3" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "45390e6114f68f718cc7a830514a96f903cccd70d02a8f6d9f643ac4ba45afaf" +dependencies = [ + "bitflags 1.3.2", + "core-foundation", + "libc", +] + +[[package]] +name = "d3d12" +version = "0.20.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "b28bfe653d79bd16c77f659305b195b82bb5ce0c0eb2a4846b82ddbd77586813" +dependencies = [ + "bitflags 2.11.0", + "libloading 0.8.9", + "winapi", +] + +[[package]] +name = "document-features" +version = "0.2.12" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "d4b8a88685455ed29a21542a33abd9cb6510b6b129abadabdcef0f4c55bc8f61" +dependencies = [ + "litrs", +] + +[[package]] +name = "equivalent" +version = "1.0.2" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "877a4ace8713b0bcf2a4e7eec82529c029f1d0619886d18145fea96c3ffe5c0f" + +[[package]] +name = "foldhash" +version = "0.1.5" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "d9c4f5dac5e15c24eb999c26181a6ca40b39fe946cbe4c263c7209467bc83af2" + +[[package]] +name = "foreign-types" +version = "0.5.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "d737d9aa519fb7b749cbc3b962edcf310a8dd1f4b67c91c4f83975dbdd17d965" +dependencies = [ + "foreign-types-macros", + "foreign-types-shared", +] + +[[package]] +name = "foreign-types-macros" +version = "0.2.3" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "1a5c6c585bc94aaf2c7b51dd4c2ba22680844aba4c687be581871a6f518c5742" +dependencies = [ + "proc-macro2", + "quote", + "syn 2.0.117", +] + +[[package]] +name = "foreign-types-shared" +version = "0.3.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "aa9a19cbb55df58761df49b23516a86d432839add4af60fc256da840f66ed35b" + +[[package]] +name = "futures-core" +version = "0.3.32" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "7e3450815272ef58cec6d564423f6e755e25379b217b0bc688e295ba24df6b1d" + +[[package]] +name = "futures-task" +version = "0.3.32" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "037711b3d59c33004d3856fbdc83b99d4ff37a24768fa1be9ce3538a1cde4393" + +[[package]] +name = "futures-util" +version = "0.3.32" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "389ca41296e6190b48053de0321d02a77f32f8a5d2461dd38762c0593805c6d6" +dependencies = [ + "futures-core", + "futures-task", + "pin-project-lite", + "slab", +] + +[[package]] +name = "gl_generator" +version = "0.14.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "1a95dfc23a2b4a9a2f5ab41d194f8bfda3cabec42af4e39f08c339eb2a0c124d" +dependencies = [ + "khronos_api", + "log", + "xml-rs", +] + +[[package]] +name = "glow" +version = "0.13.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "bd348e04c43b32574f2de31c8bb397d96c9fcfa1371bd4ca6d8bdc464ab121b1" +dependencies = [ + "js-sys", + "slotmap", + "wasm-bindgen", + "web-sys", +] + +[[package]] +name = "glutin_wgl_sys" +version = "0.5.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "6c8098adac955faa2d31079b65dc48841251f69efd3ac25477903fc424362ead" +dependencies = [ + "gl_generator", +] + +[[package]] +name = "gpu-alloc" +version = "0.6.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "fbcd2dba93594b227a1f57ee09b8b9da8892c34d55aa332e034a228d0fe6a171" +dependencies = [ + "bitflags 2.11.0", + "gpu-alloc-types", +] + +[[package]] +name = "gpu-alloc-types" +version = "0.3.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "98ff03b468aa837d70984d55f5d3f846f6ec31fe34bbb97c4f85219caeee1ca4" +dependencies = [ + "bitflags 2.11.0", +] + +[[package]] +name = "gpu-allocator" +version = "0.25.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "6f56f6318968d03c18e1bcf4857ff88c61157e9da8e47c5f29055d60e1228884" +dependencies = [ + "log", + "presser", + "thiserror", + "winapi", + "windows", +] + +[[package]] +name = "gpu-descriptor" +version = "0.3.2" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "b89c83349105e3732062a895becfc71a8f921bb71ecbbdd8ff99263e3b53a0ca" +dependencies = [ + "bitflags 2.11.0", + "gpu-descriptor-types", + "hashbrown 0.15.5", +] + +[[package]] +name = "gpu-descriptor-types" +version = "0.2.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "fdf242682df893b86f33a73828fb09ca4b2d3bb6cc95249707fc684d27484b91" +dependencies = [ + "bitflags 2.11.0", +] + +[[package]] +name = "hashbrown" +version = "0.15.5" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "9229cfe53dfd69f0609a49f65461bd93001ea1ef889cd5529dd176593f5338a1" +dependencies = [ + "foldhash", +] + +[[package]] +name = "hashbrown" +version = "0.16.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "841d1cc9bed7f9236f321df977030373f4a4163ae1a7dbfe1a51a2c1a51d9100" + +[[package]] +name = "hassle-rs" +version = "0.11.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "af2a7e73e1f34c48da31fb668a907f250794837e08faa144fd24f0b8b741e890" +dependencies = [ + "bitflags 2.11.0", + "com", + "libc", + "libloading 0.8.9", + "thiserror", + "widestring", + "winapi", +] + +[[package]] +name = "hexf-parse" +version = "0.2.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "dfa686283ad6dd069f105e5ab091b04c62850d3e4cf5d67debad1933f55023df" + +[[package]] +name = "indexmap" +version = "2.13.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "7714e70437a7dc3ac8eb7e6f8df75fd8eb422675fc7678aff7364301092b1017" +dependencies = [ + "equivalent", + "hashbrown 0.16.1", +] + +[[package]] +name = "jni-sys" +version = "0.3.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "8eaf4bc02d17cbdd7ff4c7438cafcdf7fb9a4613313ad11b4f8fefe7d3fa0130" + +[[package]] +name = "js-sys" +version = "0.3.91" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "b49715b7073f385ba4bc528e5747d02e66cb39c6146efb66b781f131f0fb399c" +dependencies = [ + "once_cell", + "wasm-bindgen", +] + +[[package]] +name = "khronos-egl" +version = "6.0.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "6aae1df220ece3c0ada96b8153459b67eebe9ae9212258bb0134ae60416fdf76" +dependencies = [ + "libc", + "libloading 0.8.9", + "pkg-config", +] + +[[package]] +name = "khronos_api" +version = "3.1.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "e2db585e1d738fc771bf08a151420d3ed193d9d895a36df7f6f8a9456b911ddc" + +[[package]] +name = "libc" +version = "0.2.183" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "b5b646652bf6661599e1da8901b3b9522896f01e736bad5f723fe7a3a27f899d" + +[[package]] +name = "libloading" +version = "0.7.4" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "b67380fd3b2fbe7527a606e18729d21c6f3951633d0500574c4dc22d2d638b9f" +dependencies = [ + "cfg-if", + "winapi", +] + +[[package]] +name = "libloading" +version = "0.8.9" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "d7c4b02199fee7c5d21a5ae7d8cfa79a6ef5bb2fc834d6e9058e89c825efdc55" +dependencies = [ + "cfg-if", + "windows-link", +] + +[[package]] +name = "litrs" +version = "1.0.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "11d3d7f243d5c5a8b9bb5d6dd2b1602c0cb0b9db1621bafc7ed66e35ff9fe092" + +[[package]] +name = "lock_api" +version = "0.4.14" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "224399e74b87b5f3557511d98dff8b14089b3dadafcab6bb93eab67d3aace965" +dependencies = [ + "scopeguard", +] + +[[package]] +name = "log" +version = "0.4.29" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "5e5032e24019045c762d3c0f28f5b6b8bbf38563a65908389bf7978758920897" + +[[package]] +name = "malloc_buf" +version = "0.0.6" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "62bb907fe88d54d8d9ce32a3cceab4218ed2f6b7d35617cafe9adf84e43919cb" +dependencies = [ + "libc", +] + +[[package]] +name = "metal" +version = "0.28.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "5637e166ea14be6063a3f8ba5ccb9a4159df7d8f6d61c02fc3d480b1f90dcfcb" +dependencies = [ + "bitflags 2.11.0", + "block", + "core-graphics-types", + "foreign-types", + "log", + "objc", + "paste", +] + +[[package]] +name = "naga" +version = "0.20.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "e536ae46fcab0876853bd4a632ede5df4b1c2527a58f6c5a4150fe86be858231" +dependencies = [ + "arrayvec", + "bit-set", + "bitflags 2.11.0", + "codespan-reporting", + "hexf-parse", + "indexmap", + "log", + "num-traits", + "rustc-hash", + "spirv", + "termcolor", + "thiserror", + "unicode-xid", +] + +[[package]] +name = "ndk-sys" +version = "0.5.0+25.2.9519653" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "8c196769dd60fd4f363e11d948139556a344e79d451aeb2fa2fd040738ef7691" +dependencies = [ + "jni-sys", +] + +[[package]] +name = "num-traits" +version = "0.2.19" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "071dfc062690e90b734c0b2273ce72ad0ffa95f0c74596bc250dcfd960262841" +dependencies = [ + "autocfg", +] + +[[package]] +name = "objc" +version = "0.2.7" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "915b1b472bc21c53464d6c8461c9d3af805ba1ef837e1cac254428f4a77177b1" +dependencies = [ + "malloc_buf", +] + +[[package]] +name = "once_cell" +version = "1.21.3" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "42f5e15c9953c5e4ccceeb2e7382a716482c34515315f7b03532b8b4e8393d2d" + +[[package]] +name = "parking_lot" +version = "0.12.5" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "93857453250e3077bd71ff98b6a65ea6621a19bb0f559a85248955ac12c45a1a" +dependencies = [ + "lock_api", + "parking_lot_core", +] + +[[package]] +name = "parking_lot_core" +version = "0.9.12" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "2621685985a2ebf1c516881c026032ac7deafcda1a2c9b7850dc81e3dfcb64c1" +dependencies = [ + "cfg-if", + "libc", + "redox_syscall", + "smallvec", + "windows-link", +] + +[[package]] +name = "paste" +version = "1.0.15" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "57c0d7b74b563b49d38dae00a0c37d4d6de9b432382b2892f0574ddcae73fd0a" + +[[package]] +name = "pin-project-lite" +version = "0.2.17" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "a89322df9ebe1c1578d689c92318e070967d1042b512afbe49518723f4e6d5cd" + +[[package]] +name = "pkg-config" +version = "0.3.32" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "7edddbd0b52d732b21ad9a5fab5c704c14cd949e5e9a1ec5929a24fded1b904c" + +[[package]] +name = "pollster" +version = "0.3.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "22686f4785f02a4fcc856d3b3bb19bf6c8160d103f7a99cc258bddd0251dc7f2" + +[[package]] +name = "presser" +version = "0.3.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "e8cf8e6a8aa66ce33f63993ffc4ea4271eb5b0530a9002db8455ea6050c77bfa" + +[[package]] +name = "proc-macro2" +version = "1.0.106" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "8fd00f0bb2e90d81d1044c2b32617f68fcb9fa3bb7640c23e9c748e53fb30934" +dependencies = [ + "unicode-ident", +] + +[[package]] +name = "profiling" +version = "1.0.17" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "3eb8486b569e12e2c32ad3e204dbaba5e4b5b216e9367044f25f1dba42341773" + +[[package]] +name = "quote" +version = "1.0.45" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "41f2619966050689382d2b44f664f4bc593e129785a36d6ee376ddf37259b924" +dependencies = [ + "proc-macro2", +] + +[[package]] +name = "range-alloc" +version = "0.1.5" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "ca45419789ae5a7899559e9512e58ca889e41f04f1f2445e9f4b290ceccd1d08" + +[[package]] +name = "raw-window-handle" +version = "0.6.2" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "20675572f6f24e9e76ef639bc5552774ed45f1c30e2951e1e99c59888861c539" + +[[package]] +name = "redox_syscall" +version = "0.5.18" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "ed2bf2547551a7053d6fdfafda3f938979645c44812fbfcda098faae3f1a362d" +dependencies = [ + "bitflags 2.11.0", +] + +[[package]] +name = "renderdoc-sys" +version = "1.1.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "19b30a45b0cd0bcca8037f3d0dc3421eaf95327a17cad11964fb8179b4fc4832" + +[[package]] +name = "rustc-hash" +version = "1.1.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "08d43f7aa6b08d49f382cde6a7982047c3426db949b1424bc4b7ec9ae12c6ce2" + +[[package]] +name = "rustversion" +version = "1.0.22" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "b39cdef0fa800fc44525c84ccb54a029961a8215f9619753635a9c0d2538d46d" + +[[package]] +name = "scopeguard" +version = "1.2.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "94143f37725109f92c262ed2cf5e59bce7498c01bcc1502d7b9afe439a4e9f49" + +[[package]] +name = "slab" +version = "0.4.12" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "0c790de23124f9ab44544d7ac05d60440adc586479ce501c1d6d7da3cd8c9cf5" + +[[package]] +name = "slotmap" +version = "1.1.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "bdd58c3c93c3d278ca835519292445cb4b0d4dc59ccfdf7ceadaab3f8aeb4038" +dependencies = [ + "version_check", +] + +[[package]] +name = "smallvec" +version = "1.15.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "67b1b7a3b5fe4f1376887184045fcf45c69e92af734b7aaddc05fb777b6fbd03" + +[[package]] +name = "sonar_wgpu" +version = "0.1.0" +dependencies = [ + "bytemuck", + "pollster", + "wgpu", +] + +[[package]] +name = "spirv" +version = "0.3.0+sdk-1.3.268.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "eda41003dc44290527a59b13432d4a0379379fa074b70174882adfbdfd917844" +dependencies = [ + "bitflags 2.11.0", +] + +[[package]] +name = "static_assertions" +version = "1.1.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "a2eb9349b6444b326872e140eb1cf5e7c522154d69e7a0ffb0fb81c06b37543f" + +[[package]] +name = "syn" +version = "1.0.109" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "72b64191b275b66ffe2469e8af2c1cfe3bafa67b529ead792a6d0160888b4237" +dependencies = [ + "proc-macro2", + "quote", + "unicode-ident", +] + +[[package]] +name = "syn" +version = "2.0.117" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "e665b8803e7b1d2a727f4023456bbbbe74da67099c585258af0ad9c5013b9b99" +dependencies = [ + "proc-macro2", + "quote", + "unicode-ident", +] + +[[package]] +name = "termcolor" +version = "1.4.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "06794f8f6c5c898b3275aebefa6b8a1cb24cd2c6c79397ab15774837a0bc5755" +dependencies = [ + "winapi-util", +] + +[[package]] +name = "thiserror" +version = "1.0.69" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "b6aaf5339b578ea85b50e080feb250a3e8ae8cfcdff9a461c9ec2904bc923f52" +dependencies = [ + "thiserror-impl", +] + +[[package]] +name = "thiserror-impl" +version = "1.0.69" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "4fee6c4efc90059e10f81e6d42c60a18f76588c3d74cb83a0b242a2b6c7504c1" +dependencies = [ + "proc-macro2", + "quote", + "syn 2.0.117", +] + +[[package]] +name = "unicode-ident" +version = "1.0.24" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "e6e4313cd5fcd3dad5cafa179702e2b244f760991f45397d14d4ebf38247da75" + +[[package]] +name = "unicode-width" +version = "0.1.14" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "7dd6e30e90baa6f72411720665d41d89b9a3d039dc45b8faea1ddd07f617f6af" + +[[package]] +name = "unicode-xid" +version = "0.2.6" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "ebc1c04c71510c7f702b52b7c350734c9ff1295c464a03335b00bb84fc54f853" + +[[package]] +name = "version_check" +version = "0.9.5" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "0b928f33d975fc6ad9f86c8f283853ad26bdd5b10b7f1542aa2fa15e2289105a" + +[[package]] +name = "wasm-bindgen" +version = "0.2.114" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "6532f9a5c1ece3798cb1c2cfdba640b9b3ba884f5db45973a6f442510a87d38e" +dependencies = [ + "cfg-if", + "once_cell", + "rustversion", + "wasm-bindgen-macro", + "wasm-bindgen-shared", +] + +[[package]] +name = "wasm-bindgen-futures" +version = "0.4.64" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "e9c5522b3a28661442748e09d40924dfb9ca614b21c00d3fd135720e48b67db8" +dependencies = [ + "cfg-if", + "futures-util", + "js-sys", + "once_cell", + "wasm-bindgen", + "web-sys", +] + +[[package]] +name = "wasm-bindgen-macro" +version = "0.2.114" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "18a2d50fcf105fb33bb15f00e7a77b772945a2ee45dcf454961fd843e74c18e6" +dependencies = [ + "quote", + "wasm-bindgen-macro-support", +] + +[[package]] +name = "wasm-bindgen-macro-support" +version = "0.2.114" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "03ce4caeaac547cdf713d280eda22a730824dd11e6b8c3ca9e42247b25c631e3" +dependencies = [ + "bumpalo", + "proc-macro2", + "quote", + "syn 2.0.117", + "wasm-bindgen-shared", +] + +[[package]] +name = "wasm-bindgen-shared" +version = "0.2.114" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "75a326b8c223ee17883a4251907455a2431acc2791c98c26279376490c378c16" +dependencies = [ + "unicode-ident", +] + +[[package]] +name = "web-sys" +version = "0.3.91" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "854ba17bb104abfb26ba36da9729addc7ce7f06f5c0f90f3c391f8461cca21f9" +dependencies = [ + "js-sys", + "wasm-bindgen", +] + +[[package]] +name = "wgpu" +version = "0.20.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "90e37c7b9921b75dfd26dd973fdcbce36f13dfa6e2dc82aece584e0ed48c355c" +dependencies = [ + "arrayvec", + "cfg-if", + "cfg_aliases", + "document-features", + "js-sys", + "log", + "naga", + "parking_lot", + "profiling", + "raw-window-handle", + "smallvec", + "static_assertions", + "wasm-bindgen", + "wasm-bindgen-futures", + "web-sys", + "wgpu-core", + "wgpu-hal", + "wgpu-types", +] + +[[package]] +name = "wgpu-core" +version = "0.21.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "d50819ab545b867d8a454d1d756b90cd5f15da1f2943334ca314af10583c9d39" +dependencies = [ + "arrayvec", + "bit-vec", + "bitflags 2.11.0", + "cfg_aliases", + "codespan-reporting", + "document-features", + "indexmap", + "log", + "naga", + "once_cell", + "parking_lot", + "profiling", + "raw-window-handle", + "rustc-hash", + "smallvec", + "thiserror", + "web-sys", + "wgpu-hal", + "wgpu-types", +] + +[[package]] +name = "wgpu-hal" +version = "0.21.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "172e490a87295564f3fcc0f165798d87386f6231b04d4548bca458cbbfd63222" +dependencies = [ + "android_system_properties", + "arrayvec", + "ash", + "bit-set", + "bitflags 2.11.0", + "block", + "cfg_aliases", + "core-graphics-types", + "d3d12", + "glow", + "glutin_wgl_sys", + "gpu-alloc", + "gpu-allocator", + "gpu-descriptor", + "hassle-rs", + "js-sys", + "khronos-egl", + "libc", + "libloading 0.8.9", + "log", + "metal", + "naga", + "ndk-sys", + "objc", + "once_cell", + "parking_lot", + "profiling", + "range-alloc", + "raw-window-handle", + "renderdoc-sys", + "rustc-hash", + "smallvec", + "thiserror", + "wasm-bindgen", + "web-sys", + "wgpu-types", + "winapi", +] + +[[package]] +name = "wgpu-types" +version = "0.20.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "1353d9a46bff7f955a680577f34c69122628cc2076e1d6f3a9be6ef00ae793ef" +dependencies = [ + "bitflags 2.11.0", + "js-sys", + "web-sys", +] + +[[package]] +name = "widestring" +version = "1.2.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "72069c3113ab32ab29e5584db3c6ec55d416895e60715417b5b883a357c3e471" + +[[package]] +name = "winapi" +version = "0.3.9" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "5c839a674fcd7a98952e593242ea400abe93992746761e38641405d28b00f419" +dependencies = [ + "winapi-i686-pc-windows-gnu", + "winapi-x86_64-pc-windows-gnu", +] + +[[package]] +name = "winapi-i686-pc-windows-gnu" +version = "0.4.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "ac3b87c63620426dd9b991e5ce0329eff545bccbbb34f3be09ff6fb6ab51b7b6" + +[[package]] +name = "winapi-util" +version = "0.1.11" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "c2a7b1c03c876122aa43f3020e6c3c3ee5c05081c9a00739faf7503aeba10d22" +dependencies = [ + "windows-sys", +] + +[[package]] +name = "winapi-x86_64-pc-windows-gnu" +version = "0.4.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "712e227841d057c1ee1cd2fb22fa7e5a5461ae8e48fa2ca79ec42cfc1931183f" + +[[package]] +name = "windows" +version = "0.52.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "e48a53791691ab099e5e2ad123536d0fff50652600abaf43bbf952894110d0be" +dependencies = [ + "windows-core", + "windows-targets", +] + +[[package]] +name = "windows-core" +version = "0.52.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "33ab640c8d7e35bf8ba19b884ba838ceb4fba93a4e8c65a9059d08afcfc683d9" +dependencies = [ + "windows-targets", +] + +[[package]] +name = "windows-link" +version = "0.2.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "f0805222e57f7521d6a62e36fa9163bc891acd422f971defe97d64e70d0a4fe5" + +[[package]] +name = "windows-sys" +version = "0.61.2" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "ae137229bcbd6cdf0f7b80a31df61766145077ddf49416a728b02cb3921ff3fc" +dependencies = [ + "windows-link", +] + +[[package]] +name = "windows-targets" +version = "0.52.6" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "9b724f72796e036ab90c1021d4780d4d3d648aca59e491e6b98e725b84e99973" +dependencies = [ + "windows_aarch64_gnullvm", + "windows_aarch64_msvc", + "windows_i686_gnu", + "windows_i686_gnullvm", + "windows_i686_msvc", + "windows_x86_64_gnu", + "windows_x86_64_gnullvm", + "windows_x86_64_msvc", +] + +[[package]] +name = "windows_aarch64_gnullvm" +version = "0.52.6" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "32a4622180e7a0ec044bb555404c800bc9fd9ec262ec147edd5989ccd0c02cd3" + +[[package]] +name = "windows_aarch64_msvc" +version = "0.52.6" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "09ec2a7bb152e2252b53fa7803150007879548bc709c039df7627cabbd05d469" + +[[package]] +name = "windows_i686_gnu" +version = "0.52.6" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "8e9b5ad5ab802e97eb8e295ac6720e509ee4c243f69d781394014ebfe8bbfa0b" + +[[package]] +name = "windows_i686_gnullvm" +version = "0.52.6" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "0eee52d38c090b3caa76c563b86c3a4bd71ef1a819287c19d586d7334ae8ed66" + +[[package]] +name = "windows_i686_msvc" +version = "0.52.6" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "240948bc05c5e7c6dabba28bf89d89ffce3e303022809e73deaefe4f6ec56c66" + +[[package]] +name = "windows_x86_64_gnu" +version = "0.52.6" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "147a5c80aabfbf0c7d901cb5895d1de30ef2907eb21fbbab29ca94c5b08b1a78" + +[[package]] +name = "windows_x86_64_gnullvm" +version = "0.52.6" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "24d5b23dc417412679681396f2b49f3de8c1473deb516bd34410872eff51ed0d" + +[[package]] +name = "windows_x86_64_msvc" +version = "0.52.6" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "589f6da84c646204747d1270a2a5661ea66ed1cced2631d546fdfb155959f9ec" + +[[package]] +name = "xml-rs" +version = "0.8.28" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "3ae8337f8a065cfc972643663ea4279e04e7256de865aa66fe25cec5fb912d3f" diff --git a/gazebo/posim_gz_multibeam_sonar/wgpu_vendor/sonar_wgpu_rust/Cargo.toml b/gazebo/posim_gz_multibeam_sonar/wgpu_vendor/sonar_wgpu_rust/Cargo.toml new file mode 100644 index 00000000..78965cd2 --- /dev/null +++ b/gazebo/posim_gz_multibeam_sonar/wgpu_vendor/sonar_wgpu_rust/Cargo.toml @@ -0,0 +1,16 @@ +[package] +name = "sonar_wgpu" +version = "0.1.0" +edition = "2021" + +[lib] +crate-type = ["staticlib"] + +[dependencies] +wgpu = { version = "0.20" } +bytemuck = { version = "1", features = ["derive"] } +pollster = "0.3" + +[profile.release] +opt-level = 3 +lto = true diff --git a/gazebo/posim_gz_multibeam_sonar/wgpu_vendor/sonar_wgpu_rust/src/lib.rs b/gazebo/posim_gz_multibeam_sonar/wgpu_vendor/sonar_wgpu_rust/src/lib.rs new file mode 100644 index 00000000..a8c6ea19 --- /dev/null +++ b/gazebo/posim_gz_multibeam_sonar/wgpu_vendor/sonar_wgpu_rust/src/lib.rs @@ -0,0 +1,534 @@ +//! C-FFI interface, Buffer initialization and write ups and reads + Bind groups created + encoder setup + memory free +//! + Compute passes for each shader + +mod pipeline; // GPU context + buffer management + +use bytemuck::{Pod, Zeroable}; +use std::fs::OpenOptions; +use std::io::Write; +use std::sync::atomic::{AtomicU64, Ordering}; +use wgpu::util::DeviceExt; + +#[repr(C)] +#[derive(Clone, Copy, Pod, Zeroable)] +struct BackscatterParams { + n_beams: u32, + n_rays: u32, + n_freq: u32, + ray_skips: u32, + sound_speed: f32, + max_distance: f32, + source_level: f32, + attenuation: f32, + sensor_gain: f32, + bandwidth: f32, + seed_lo: u32, + seed_hi: u32, + frame_lo: u32, // WGSL has no native u64 type. + frame_hi: u32, // Split frame index into two u32 - frame_lo and frame_hi + area_scaler: f32, + h_fov: f32, + v_fov: f32, + _pad0: u32, + _pad1: u32, + _pad2: u32, +} + +#[repr(C)] +#[derive(Clone, Copy, Pod, Zeroable)] +struct MatmulParams { + n_beams: u32, + n_freq: u32, + beam_corrector_sum: f32, +} + +#[repr(C)] +#[derive(Clone, Copy, Pod, Zeroable)] +struct FftParams { + n_beams: u32, + n_freq: u32, // actual number of frequency bins (may not be power-of-2) + padded_n: u32, // next power of 2 >= n_freq (actual FFT size, <= 4096) + log2_n: u32, // log2(padded_n) +} + +fn emit_debug_timing(line: &str) { + eprintln!("{line}"); + + if let Ok(mut file) = OpenOptions::new() + .create(true) + .append(true) + .open("wgpu_debug_timings.txt") + { + let _ = writeln!(file, "{line}"); + let _ = file.flush(); + } +} + +/// The function receives flat arrays from C++ and returns a heap-allocated array. +/// +/// # Safety +/// +/// Each input pointer must be null or valid for reads of its full length: +/// `depth_flat` and `reflectivity_flat` `n_beams * n_rays` floats, `normal_flat` +/// `n_beams * n_rays * 3`, `beam_corrector_flat` `n_beams * n_beams` and `window_flat` `n_freq`. +/// A non-null result holds `n_beams * n_freq * 2` floats and must be released with +/// [`sonar_wgpu_free`]. +#[no_mangle] +pub unsafe extern "C" fn sonar_wgpu_compute( + depth_flat: *const f32, + normal_flat: *const f32, + reflectivity_flat: *const f32, + beam_corrector_flat: *const f32, + window_flat: *const f32, + n_beams: u32, + n_rays: u32, + n_freq: u32, + ray_skips: u32, + sound_speed: f32, + max_distance: f32, + source_level: f32, + attenuation: f32, + sensor_gain: f32, + bandwidth: f32, + beam_corrector_sum: f32, + area_scaler: f32, + h_fov: f32, + v_fov: f32, + frame_index: u64, + seed: u64, + debug_flag: bool, +) -> *mut f32 { + if depth_flat.is_null() + || normal_flat.is_null() + || reflectivity_flat.is_null() + || beam_corrector_flat.is_null() + || window_flat.is_null() + { + return std::ptr::null_mut(); + } + + let n_beams_us = n_beams as usize; + let n_rays_us = n_rays as usize; + let n_freq_us = n_freq as usize; + + if n_beams_us == 0 || n_rays_us == 0 || n_freq_us == 0 { + return std::ptr::null_mut(); + } + + use std::sync::Once; + + let depth_len = n_beams_us * n_rays_us; + let normal_len = depth_len * 3; + let refl_len = depth_len; + let bc_len = n_beams_us * n_beams_us; + let spectrum_len = n_beams_us * n_freq_us; + + let depth = unsafe { std::slice::from_raw_parts(depth_flat, depth_len) }; + let normals = unsafe { std::slice::from_raw_parts(normal_flat, normal_len) }; + let reflectivity = unsafe { std::slice::from_raw_parts(reflectivity_flat, refl_len) }; + let beam_corrector = unsafe { std::slice::from_raw_parts(beam_corrector_flat, bc_len) }; + + let ctx = match pipeline::get_or_init() { + Some(c) => { + static GPU_LOG: Once = Once::new(); + GPU_LOG.call_once(|| eprintln!("[sonar_wgpu] GPU device acquired -> running on GPU.")); + c + } + None => { + static FALLBACK_LOG: Once = Once::new(); + FALLBACK_LOG.call_once(|| eprintln!("[sonar_wgpu] GPU init failed -> returning null.")); + return std::ptr::null_mut(); + } + }; + let device = &ctx.device; + let queue = &ctx.queue; + + static GPU_FRAME_COUNT: AtomicU64 = AtomicU64::new(0); + let gpu_frame = GPU_FRAME_COUNT.fetch_add(1, Ordering::Relaxed); + let gpu_t0 = std::time::Instant::now(); + + let window_vals = unsafe { std::slice::from_raw_parts(window_flat, n_freq_us) }; + + // ---- Per-frame uniform buffers (tiny, always recreated) ---- + let backscatter_params = BackscatterParams { + n_beams, + n_rays, + n_freq, + ray_skips, + sound_speed, + max_distance, + source_level, + attenuation, + sensor_gain, + bandwidth, + seed_lo: (seed & 0xFFFF_FFFF) as u32, + seed_hi: (seed >> 32) as u32, + frame_lo: (frame_index & 0xFFFF_FFFF) as u32, + frame_hi: (frame_index >> 32) as u32, + area_scaler, + h_fov, + v_fov, + _pad0: 0, + _pad1: 0, + _pad2: 0, + }; + let bs_param_buf = device.create_buffer_init(&wgpu::util::BufferInitDescriptor { + label: Some("bs_param"), + contents: bytemuck::bytes_of(&backscatter_params), + usage: wgpu::BufferUsages::UNIFORM | wgpu::BufferUsages::COPY_DST, + }); + let mm_params = MatmulParams { + n_beams, + n_freq, + beam_corrector_sum, + }; + let mm_param_buf = device.create_buffer_init(&wgpu::util::BufferInitDescriptor { + label: Some("mm_param"), + contents: bytemuck::bytes_of(&mm_params), + usage: wgpu::BufferUsages::UNIFORM | wgpu::BufferUsages::COPY_DST, + }); + + // REMOVE THIS, UNNCESSARY + // ---- Persistent data buffers: allocate once per dimension set, reuse every frame ---- + // The lock is held for the whole frame, so sonars with equal dimensions can share a set. + let mut buf_guard = ctx.buffers.lock().unwrap(); + let key = (n_beams, n_rays, n_freq); + if !buf_guard.contains_key(&key) { + if buf_guard.len() >= pipeline::MAX_BUFFER_SETS { + buf_guard.clear(); + } + buf_guard.insert( + key, + pipeline::SonarBuffers::new(device, n_beams, n_rays, n_freq), + ); + eprintln!( + "[sonar_wgpu] Persistent GPU buffers allocated for {}×{}×{}", + n_beams, n_rays, n_freq + ); + } + let buf = &buf_guard[&key]; + + // Upload frame data into persistent buffers (DMA, no GPU alloc) + queue.write_buffer(&buf.depth_buf, 0, bytemuck::cast_slice(depth)); + queue.write_buffer(&buf.normal_buf, 0, bytemuck::cast_slice(normals)); + queue.write_buffer(&buf.refl_buf, 0, bytemuck::cast_slice(reflectivity)); + queue.write_buffer(&buf.window_buf, 0, bytemuck::cast_slice(window_vals)); + queue.write_buffer(&buf.bc_buf, 0, bytemuck::cast_slice(beam_corrector)); + // Zero atomic accumulators (reused each frame) + let zeros_i32 = vec![0i32; spectrum_len]; + queue.write_buffer(&buf.out_re_i32, 0, bytemuck::cast_slice(&zeros_i32)); + queue.write_buffer(&buf.out_im_i32, 0, bytemuck::cast_slice(&zeros_i32)); + + // ---- Bind groups (metadata only -> no GPU allocation) ---- + let bs_bg = device.create_bind_group(&wgpu::BindGroupDescriptor { + label: Some("bs_bg"), + layout: &ctx.bs_bgl, + entries: &[ + wgpu::BindGroupEntry { + binding: 0, + resource: bs_param_buf.as_entire_binding(), + }, + wgpu::BindGroupEntry { + binding: 1, + resource: buf.depth_buf.as_entire_binding(), + }, + wgpu::BindGroupEntry { + binding: 2, + resource: buf.normal_buf.as_entire_binding(), + }, + wgpu::BindGroupEntry { + binding: 3, + resource: buf.refl_buf.as_entire_binding(), + }, + wgpu::BindGroupEntry { + binding: 4, + resource: buf.out_re_i32.as_entire_binding(), + }, + wgpu::BindGroupEntry { + binding: 5, + resource: buf.out_im_i32.as_entire_binding(), + }, + wgpu::BindGroupEntry { + binding: 6, + resource: buf.window_buf.as_entire_binding(), + }, + ], + }); + let conv_bg = device.create_bind_group(&wgpu::BindGroupDescriptor { + label: Some("conv_bg"), + layout: &ctx.convert_bgl, + entries: &[ + wgpu::BindGroupEntry { + binding: 0, + resource: buf.out_re_i32.as_entire_binding(), + }, + wgpu::BindGroupEntry { + binding: 1, + resource: buf.out_im_i32.as_entire_binding(), + }, + wgpu::BindGroupEntry { + binding: 2, + resource: buf.mm_re_in.as_entire_binding(), + }, + wgpu::BindGroupEntry { + binding: 3, + resource: buf.mm_im_in.as_entire_binding(), + }, + ], + }); + let mm_re_bg = device.create_bind_group(&wgpu::BindGroupDescriptor { + label: Some("mm_re_bg"), + layout: &ctx.mm_bgl, + entries: &[ + wgpu::BindGroupEntry { + binding: 0, + resource: mm_param_buf.as_entire_binding(), + }, + wgpu::BindGroupEntry { + binding: 1, + resource: buf.bc_buf.as_entire_binding(), + }, + wgpu::BindGroupEntry { + binding: 2, + resource: buf.mm_re_in.as_entire_binding(), + }, + wgpu::BindGroupEntry { + binding: 3, + resource: buf.mm_re_out.as_entire_binding(), + }, + ], + }); + let mm_im_bg = device.create_bind_group(&wgpu::BindGroupDescriptor { + label: Some("mm_im_bg"), + layout: &ctx.mm_bgl, + entries: &[ + wgpu::BindGroupEntry { + binding: 0, + resource: mm_param_buf.as_entire_binding(), + }, + wgpu::BindGroupEntry { + binding: 1, + resource: buf.bc_buf.as_entire_binding(), + }, + wgpu::BindGroupEntry { + binding: 2, + resource: buf.mm_im_in.as_entire_binding(), + }, + wgpu::BindGroupEntry { + binding: 3, + resource: buf.mm_im_out.as_entire_binding(), + }, + ], + }); + + let spectrum_bytes = (spectrum_len * std::mem::size_of::()) as u64; + // Zero-pad n_freq to the next power of 2 for the GPU FFT (max 4096 = shared memory limit). + let padded_n = (n_freq_us.next_power_of_two() as u32).min(4096); + + // ---- Single command encoder: all passes + staging copy in one GPU submission ---- + let mut enc = device.create_command_encoder(&wgpu::CommandEncoderDescriptor { + label: Some("sonar_enc"), + }); + + // Backscatter dispatch: depth/normals/refl → i32 fixed-point spectrum + { + let mut pass = enc.begin_compute_pass(&wgpu::ComputePassDescriptor { + label: Some("bs_pass"), + timestamp_writes: None, + }); + pass.set_pipeline(&ctx.bs_pipeline); + pass.set_bind_group(0, &bs_bg, &[]); + let gx = n_beams.div_ceil(8); + let ray_step = std::cmp::max(1, ray_skips); + let reduced_rays = n_rays.div_ceil(ray_step); + let gy = reduced_rays.div_ceil(8); + pass.dispatch_workgroups(gx, gy, 1); + } + + // Convert dispatch: i32 fixed-point → f32 (GPU-side, eliminates CPU readback round-trip) + { + let mut pass = enc.begin_compute_pass(&wgpu::ComputePassDescriptor { + label: Some("conv_pass"), + timestamp_writes: None, + }); + pass.set_pipeline(&ctx.convert_pipeline); + pass.set_bind_group(0, &conv_bg, &[]); + let gx = (spectrum_len as u32).div_ceil(64); + pass.dispatch_workgroups(gx, 1, 1); + } + + // Matmul dispatch: beam-corrector matrix multiply, re and im in same encoder + { + let gx = n_freq.div_ceil(16); + let gy = n_beams.div_ceil(16); + { + let mut pass = enc.begin_compute_pass(&wgpu::ComputePassDescriptor { + label: Some("mm_re_pass"), + timestamp_writes: None, + }); + pass.set_pipeline(&ctx.mm_pipeline); + pass.set_bind_group(0, &mm_re_bg, &[]); + pass.dispatch_workgroups(gx, gy, 1); + } + { + let mut pass = enc.begin_compute_pass(&wgpu::ComputePassDescriptor { + label: Some("mm_im_pass"), + timestamp_writes: None, + }); + pass.set_pipeline(&ctx.mm_pipeline); + pass.set_bind_group(0, &mm_im_bg, &[]); + pass.dispatch_workgroups(gx, gy, 1); + } + } + + // FFT Dispatch: always runs on GPU via zero-padding to the next power of 2 (max 4096). + // Copy MM output -> FFT working buffers -> dispatch FFT shader -> copy to staging. + enc.copy_buffer_to_buffer(&buf.mm_re_out, 0, &buf.p_re_buf, 0, spectrum_bytes); + enc.copy_buffer_to_buffer(&buf.mm_im_out, 0, &buf.p_im_buf, 0, spectrum_bytes); + + let fft_params = FftParams { + n_beams, + n_freq, + padded_n, + log2_n: padded_n.ilog2(), + }; + let fft_param_buf = device.create_buffer_init(&wgpu::util::BufferInitDescriptor { + label: Some("fft_param"), + contents: bytemuck::bytes_of(&fft_params), + usage: wgpu::BufferUsages::UNIFORM | wgpu::BufferUsages::COPY_DST, + }); + let fft_bg = device.create_bind_group(&wgpu::BindGroupDescriptor { + label: Some("fft_bg"), + layout: &ctx.fft_bgl, + entries: &[ + wgpu::BindGroupEntry { + binding: 0, + resource: fft_param_buf.as_entire_binding(), + }, + wgpu::BindGroupEntry { + binding: 1, + resource: buf.p_re_buf.as_entire_binding(), + }, + wgpu::BindGroupEntry { + binding: 2, + resource: buf.p_im_buf.as_entire_binding(), + }, + ], + }); + { + let mut pass = enc.begin_compute_pass(&wgpu::ComputePassDescriptor { + label: Some("fft_pass"), + timestamp_writes: None, + }); + pass.set_pipeline(&ctx.fft_pipeline); + pass.set_bind_group(0, &fft_bg, &[]); + pass.dispatch_workgroups(n_beams, 1, 1); + } + // Copy FFT output to staging for CPU readback. + enc.copy_buffer_to_buffer(&buf.p_re_buf, 0, &buf.stg_re, 0, spectrum_bytes); + enc.copy_buffer_to_buffer(&buf.p_im_buf, 0, &buf.stg_im, 0, spectrum_bytes); + + // Submit the entire frame in a single GPU command -> POLL 1 + queue.submit(std::iter::once(enc.finish())); + device.poll(wgpu::Maintain::Wait); + + // map_async is asynchronous -> it requests CPU access to the buffer but doesn't block. + // Map staging buffers (GPU work already done above) -> POLL 2 + let re_ok = std::sync::Arc::new(std::sync::atomic::AtomicBool::new(false)); + let im_ok = std::sync::Arc::new(std::sync::atomic::AtomicBool::new(false)); + let re_ok_cb = std::sync::Arc::clone(&re_ok); + let im_ok_cb = std::sync::Arc::clone(&im_ok); + + buf.stg_re + .slice(..) + .map_async(wgpu::MapMode::Read, move |r| { + re_ok_cb.store(r.is_ok(), std::sync::atomic::Ordering::Release); + }); + buf.stg_im + .slice(..) + .map_async(wgpu::MapMode::Read, move |r| { + im_ok_cb.store(r.is_ok(), std::sync::atomic::Ordering::Release); + }); + device.poll(wgpu::Maintain::Wait); + + let re_mapped = re_ok.load(std::sync::atomic::Ordering::Acquire); + let im_mapped = im_ok.load(std::sync::atomic::Ordering::Acquire); + if !re_mapped || !im_mapped { + eprintln!( + "[sonar_wgpu] map_async failed: stg_re_ok={}, stg_im_ok={}", + re_mapped, im_mapped + ); + buf.stg_re.unmap(); + buf.stg_im.unmap(); + return std::ptr::null_mut(); + } + + let (mut p_re, mut p_im) = { + let re_data = buf.stg_re.slice(..).get_mapped_range(); + let im_data = buf.stg_im.slice(..).get_mapped_range(); + let re = bytemuck::cast_slice::(&re_data).to_vec(); + let im = bytemuck::cast_slice::(&im_data).to_vec(); + (re, im) + }; + buf.stg_re.unmap(); + buf.stg_im.unmap(); + + // Matches CUDA's post-FFT * delta_f scaling + let delta_f = bandwidth / (n_freq as f32); + for i in 0..spectrum_len { + p_re[i] *= delta_f; + p_im[i] *= delta_f; + } + + // Pack into interleaved [re0, im0, re1, im1, ...] for C++ caller + let mut out = vec![0.0f32; spectrum_len * 2]; + for b in 0..n_beams_us { + for f in 0..n_freq_us { + let idx = b * n_freq_us + f; + out[idx * 2] = p_re[idx]; + out[idx * 2 + 1] = p_im[idx]; + } + } + + let elapsed_ms = gpu_t0.elapsed().as_secs_f64() * 1000.0; + if debug_flag { + emit_debug_timing(&format!( + "Total WGPU Calculation Wrapper Time: {:.3} ms", + elapsed_ms + )); + } + + // Log on the very first frame and then every 50 frames. + if gpu_frame == 0 || gpu_frame % 50 == 49 { + let active_rays = n_rays.div_ceil(std::cmp::max(1, ray_skips)); + eprintln!( + "[sonar_wgpu] GPU #{:<5} | {:6.1} ms | {} beams × {} rays × {} freq", + gpu_frame + 1, + elapsed_ms, + n_beams, + active_rays, + n_freq + ); + } + + let mut boxed = out.into_boxed_slice(); + let ptr = boxed.as_mut_ptr(); + std::mem::forget(boxed); + ptr +} + +/// Frees an array returned by [`sonar_wgpu_compute`]. +/// +/// # Safety +/// +/// `ptr` must be null or a pointer returned by `sonar_wgpu_compute` that has not been freed yet, +/// and `len` must be the length of that array (`n_beams * n_freq * 2`). +#[no_mangle] +pub unsafe extern "C" fn sonar_wgpu_free(ptr: *mut f32, len: usize) { + if ptr.is_null() || len == 0 { + return; + } + unsafe { + let _ = Vec::from_raw_parts(ptr, len, len); + } +} diff --git a/gazebo/posim_gz_multibeam_sonar/wgpu_vendor/sonar_wgpu_rust/src/pipeline.rs b/gazebo/posim_gz_multibeam_sonar/wgpu_vendor/sonar_wgpu_rust/src/pipeline.rs new file mode 100644 index 00000000..a536927b --- /dev/null +++ b/gazebo/posim_gz_multibeam_sonar/wgpu_vendor/sonar_wgpu_rust/src/pipeline.rs @@ -0,0 +1,446 @@ +/// Adapters, device, queue setups + BUFFERS SETUP + PIPELINE SETUPS + PIPELINE COMPILE + GPU CONTEXT INIT all in pipeline.rs +/// + Shaders compiled +use std::collections::HashMap; +use std::sync::OnceLock; + +/// Buffer-set key: (n_beams, n_rays, n_freq). +pub type BufferKey = (u32, u32, u32); + +/// Upper bound on cached buffer sets; the cache is cleared when it is full. +pub const MAX_BUFFER_SETS: usize = 8; + +/// GPU buffer lifetime: persistent allocation across frames via queue.write_buffer. +/// One set per sonar dimensions (see `GpuContext::buffers`). +pub struct SonarBuffers { + // Input data (written via queue.write_buffer each frame) + pub depth_buf: wgpu::Buffer, + pub normal_buf: wgpu::Buffer, + pub refl_buf: wgpu::Buffer, + pub window_buf: wgpu::Buffer, + pub bc_buf: wgpu::Buffer, + // Backscatter atomic i32 accumulators (zeroed via write_buffer each frame) + pub out_re_i32: wgpu::Buffer, + pub out_im_i32: wgpu::Buffer, + // Convert output / matmul input (f32) + pub mm_re_in: wgpu::Buffer, + pub mm_im_in: wgpu::Buffer, + // Matmul output / FFT input + pub mm_re_out: wgpu::Buffer, + pub mm_im_out: wgpu::Buffer, + // FFT in-place buffers + pub p_re_buf: wgpu::Buffer, + pub p_im_buf: wgpu::Buffer, + // Host-visible staging buffers for final readback + pub stg_re: wgpu::Buffer, + pub stg_im: wgpu::Buffer, +} + +impl SonarBuffers { + /// Buffer allocation for 4 pipeline stages: backscatter, convert, matmul, FFT. + /// Chain: backscatter→convert→matmul→FFT→readback. + pub fn new(device: &wgpu::Device, n_beams: u32, n_rays: u32, n_freq: u32) -> Self { + let depth_ray = (n_beams as usize) * (n_rays as usize); + let spectrum = (n_beams as usize) * (n_freq as usize); + let bc = (n_beams as usize) * (n_beams as usize); + let f32_bytes = |n: usize| (n * std::mem::size_of::()) as u64; + let i32_bytes = |n: usize| (n * std::mem::size_of::()) as u64; + + SonarBuffers { + depth_buf: device.create_buffer(&wgpu::BufferDescriptor { + label: Some("pb_depth"), + size: f32_bytes(depth_ray), + usage: wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_DST, + mapped_at_creation: false, + }), + normal_buf: device.create_buffer(&wgpu::BufferDescriptor { + label: Some("pb_normal"), + size: f32_bytes(depth_ray * 3), + usage: wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_DST, + mapped_at_creation: false, + }), + refl_buf: device.create_buffer(&wgpu::BufferDescriptor { + label: Some("pb_refl"), + size: f32_bytes(depth_ray), + usage: wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_DST, + mapped_at_creation: false, + }), + window_buf: device.create_buffer(&wgpu::BufferDescriptor { + label: Some("pb_window"), + size: f32_bytes(n_freq as usize), + usage: wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_DST, + mapped_at_creation: false, + }), + bc_buf: device.create_buffer(&wgpu::BufferDescriptor { + label: Some("pb_bc"), + size: f32_bytes(bc), + usage: wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_DST, + mapped_at_creation: false, + }), + out_re_i32: device.create_buffer(&wgpu::BufferDescriptor { + label: Some("pb_re_i32"), + size: i32_bytes(spectrum), + usage: wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_DST, + mapped_at_creation: false, + }), + out_im_i32: device.create_buffer(&wgpu::BufferDescriptor { + label: Some("pb_im_i32"), + size: i32_bytes(spectrum), + usage: wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_DST, + mapped_at_creation: false, + }), + mm_re_in: device.create_buffer(&wgpu::BufferDescriptor { + label: Some("pb_mm_re_in"), + size: f32_bytes(spectrum), + usage: wgpu::BufferUsages::STORAGE, + mapped_at_creation: false, + }), + mm_im_in: device.create_buffer(&wgpu::BufferDescriptor { + label: Some("pb_mm_im_in"), + size: f32_bytes(spectrum), + usage: wgpu::BufferUsages::STORAGE, + mapped_at_creation: false, + }), + mm_re_out: device.create_buffer(&wgpu::BufferDescriptor { + label: Some("pb_mm_re_out"), + size: f32_bytes(spectrum), + usage: wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_SRC, + mapped_at_creation: false, + }), + mm_im_out: device.create_buffer(&wgpu::BufferDescriptor { + label: Some("pb_mm_im_out"), + size: f32_bytes(spectrum), + usage: wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_SRC, + mapped_at_creation: false, + }), + p_re_buf: device.create_buffer(&wgpu::BufferDescriptor { + label: Some("pb_p_re"), + size: f32_bytes(spectrum), + usage: wgpu::BufferUsages::STORAGE + | wgpu::BufferUsages::COPY_SRC + | wgpu::BufferUsages::COPY_DST, + mapped_at_creation: false, + }), + p_im_buf: device.create_buffer(&wgpu::BufferDescriptor { + label: Some("pb_p_im"), + size: f32_bytes(spectrum), + usage: wgpu::BufferUsages::STORAGE + | wgpu::BufferUsages::COPY_SRC + | wgpu::BufferUsages::COPY_DST, + mapped_at_creation: false, + }), + stg_re: device.create_buffer(&wgpu::BufferDescriptor { + label: Some("pb_stg_re"), + size: f32_bytes(spectrum), + usage: wgpu::BufferUsages::MAP_READ | wgpu::BufferUsages::COPY_DST, + mapped_at_creation: false, + }), + stg_im: device.create_buffer(&wgpu::BufferDescriptor { + label: Some("pb_stg_im"), + size: f32_bytes(spectrum), + usage: wgpu::BufferUsages::MAP_READ | wgpu::BufferUsages::COPY_DST, + mapped_at_creation: false, + }), + } + } +} + +/// GPU compute context: device handle + compiled shader pipelines. +/// Four pipelines compiled once at init: backscatter, convert, matmul, FFT. +pub struct GpuContext { + pub device: wgpu::Device, // the GPU -> used to create everything + pub queue: wgpu::Queue, // the command queue -> used to submit work + + // Backscatter kernel: ray-surface interactions with Lambert scatter. + pub bs_pipeline: wgpu::ComputePipeline, // compiled backscatter shader + pub bs_bgl: wgpu::BindGroupLayout, // describes bs shader's buffer slots + + // i32→f32 conversion: fixed-point atomic backscatter output to float spectrum. + pub convert_pipeline: wgpu::ComputePipeline, + pub convert_bgl: wgpu::BindGroupLayout, + + // Beam correction GEMM: tiled matrix multiply for spectrum. + pub mm_pipeline: wgpu::ComputePipeline, + pub mm_bgl: wgpu::BindGroupLayout, + + // FFT kernel: power-of-2 only, N≤4096 (4096 is the shared memory limit). CPU fallback via Bluestein for arbitrary N. + pub fft_pipeline: wgpu::ComputePipeline, + pub fft_bgl: wgpu::BindGroupLayout, + + // Persistent GPU buffers across frames, one set per (n_beams, n_rays, n_freq), allocated on the + // first compute call with those dimensions. Keying by dimensions lets several sonars of + // different sizes share the context without reallocating every frame. + pub buffers: std::sync::Mutex>, +} + +static GPU_CONTEXT: OnceLock> = OnceLock::new(); + +/// Bind group layout entry: storage buffer descriptor. +fn make_storage_entry(binding: u32, read_only: bool) -> wgpu::BindGroupLayoutEntry { + wgpu::BindGroupLayoutEntry { + binding, + visibility: wgpu::ShaderStages::COMPUTE, + ty: wgpu::BindingType::Buffer { + ty: wgpu::BufferBindingType::Storage { read_only }, + has_dynamic_offset: false, + min_binding_size: None, + }, + count: None, + } +} + +/// Uniform entry: shader parameter buffer. +fn make_uniform_entry(binding: u32) -> wgpu::BindGroupLayoutEntry { + wgpu::BindGroupLayoutEntry { + binding, + visibility: wgpu::ShaderStages::COMPUTE, + ty: wgpu::BindingType::Buffer { + ty: wgpu::BufferBindingType::Uniform, + has_dynamic_offset: false, + min_binding_size: None, + }, + count: None, + } +} + +/// Bind group layout builder: defines buffer bindings for each shader. +fn build_bgls( + device: &wgpu::Device, +) -> ( + wgpu::BindGroupLayout, + wgpu::BindGroupLayout, + wgpu::BindGroupLayout, + wgpu::BindGroupLayout, +) { + // Backscatter kernel params: uniform params, depth, normals, reflectivity, output accumulators, window. + let bs_bgl = device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor { + label: Some("bs_bgl"), + entries: &[ + make_uniform_entry(0), // (params) + make_storage_entry(1, true), // (depth) true = read_only + make_storage_entry(2, true), // (normal) + make_storage_entry(3, true), // (reflectivity) + make_storage_entry(4, false), // (out_re_i32) false = read_write + make_storage_entry(5, false), // (out_im_i32) + make_storage_entry(6, true), // (window) + ], + }); + + // Matmul (GEMM) kernel params: uniform params, beam_corrector, spectrum_in, spectrum_out. + let mm_bgl = device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor { + label: Some("mm_bgl"), + entries: &[ + make_uniform_entry(0), + make_storage_entry(1, true), + make_storage_entry(2, true), + make_storage_entry(3, false), + ], + }); + + // FFT kernel params: uniform params, spectrum_re, spectrum_im (in-place). + let fft_bgl = device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor { + label: Some("fft_bgl"), + entries: &[ + make_uniform_entry(0), + make_storage_entry(1, false), + make_storage_entry(2, false), + ], + }); + + // Convert kernel params: int32 input (re/im), float32 output (re/im). + let convert_bgl = device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor { + label: Some("convert_bgl"), + entries: &[ + make_storage_entry(0, true), + make_storage_entry(1, true), + make_storage_entry(2, false), + make_storage_entry(3, false), + ], + }); + + (bs_bgl, mm_bgl, fft_bgl, convert_bgl) +} + +/// Pipeline compilation: shader module compilation + pipeline layout binding (one-time init). +fn compile_pipeline( + device: &wgpu::Device, + label: &str, + wgsl: &str, + bgl: &wgpu::BindGroupLayout, +) -> wgpu::ComputePipeline { + let shader = device.create_shader_module(wgpu::ShaderModuleDescriptor { + label: Some(label), + source: wgpu::ShaderSource::Wgsl(wgsl.into()), + }); + let layout = device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor { + label: Some(label), + bind_group_layouts: &[bgl], + push_constant_ranges: &[], + }); + device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor { + label: Some(label), + layout: Some(&layout), + module: &shader, + entry_point: "main", + compilation_options: wgpu::PipelineCompilationOptions::default(), + }) +} + +// Default backend sets, tried in order. WGPU_BACKEND replaces them with a single set. +// Linux tries Vulkan only: the GL backend's EGL setup panics inside Gazebo. +#[cfg(target_os = "linux")] +const BACKEND_ATTEMPTS: &[(wgpu::Backends, &str)] = &[(wgpu::Backends::VULKAN, "vulkan")]; + +#[cfg(target_os = "windows")] +const BACKEND_ATTEMPTS: &[(wgpu::Backends, &str)] = &[ + (wgpu::Backends::DX12, "dx12"), + (wgpu::Backends::all(), "all"), +]; + +#[cfg(target_os = "macos")] +const BACKEND_ATTEMPTS: &[(wgpu::Backends, &str)] = &[ + (wgpu::Backends::METAL, "metal"), + (wgpu::Backends::all(), "all"), +]; + +#[cfg(not(any(target_os = "linux", target_os = "windows", target_os = "macos")))] +const BACKEND_ATTEMPTS: &[(wgpu::Backends, &str)] = &[(wgpu::Backends::all(), "all")]; + +/// Pick an adapter. WGPU_ADAPTER_NAME selects the first adapter whose name contains it +/// (case-insensitive); otherwise wgpu chooses by WGPU_POWER_PREF (default: high performance). +fn select_adapter( + instance: &wgpu::Instance, + backends: wgpu::Backends, + label: &str, +) -> Option { + if let Ok(wanted) = std::env::var("WGPU_ADAPTER_NAME") { + let wanted = wanted.to_lowercase(); + let adapters = instance.enumerate_adapters(backends); + let names: Vec = adapters.iter().map(|a| a.get_info().name).collect(); + let found = adapters + .into_iter() + .find(|a| a.get_info().name.to_lowercase().contains(&wanted)); + if found.is_none() { + eprintln!( + "[sonar_wgpu] [{label}] no adapter matches WGPU_ADAPTER_NAME='{wanted}'; available: {names:?}" + ); + } + return found; + } + + pollster::block_on( + instance.request_adapter(&wgpu::RequestAdapterOptions { + power_preference: wgpu::util::power_preference_from_env() + .unwrap_or(wgpu::PowerPreference::HighPerformance), + compatible_surface: None, + force_fallback_adapter: false, + }), + ) +} + +fn try_init_gpu_context(backends: wgpu::Backends, label: &str) -> Option { + eprintln!("[sonar_wgpu] Trying backend set: {label}"); + + // Catch panic to prevent GPU init failure from crashing Gazebo.Whtat + let result = std::panic::catch_unwind(|| { + let instance = wgpu::Instance::new(wgpu::InstanceDescriptor { + backends, + ..Default::default() + }); + + let adapter = select_adapter(&instance, backends, label)?; + + eprintln!( + "[sonar_wgpu] [{label}] selected adapter: {}", + adapter.get_info().name + ); + let (device, queue) = pollster::block_on(adapter.request_device( + &wgpu::DeviceDescriptor { + label: Some("sonar_wgpu_device"), + required_features: wgpu::Features::empty(), + required_limits: wgpu::Limits::default(), + }, + None, + )) + .map_err(|e| { + eprintln!("[sonar_wgpu] [{label}] request_device failed: {e}"); + e + }) + .ok()?; + + // One-time shader compilation. + let t0 = std::time::Instant::now(); + let (bs_bgl, mm_bgl, fft_bgl, convert_bgl) = build_bgls(&device); + let bs_pipeline = compile_pipeline( + &device, + "bs_pipeline", + include_str!("shaders/backscatter.wgsl"), + &bs_bgl, + ); + let convert_pipeline = compile_pipeline( + &device, + "convert_pipeline", + include_str!("shaders/convert.wgsl"), + &convert_bgl, + ); + let mm_pipeline = compile_pipeline( + &device, + "mm_pipeline", + include_str!("shaders/matmul.wgsl"), + &mm_bgl, + ); + let fft_pipeline = compile_pipeline( + &device, + "fft_pipeline", + include_str!("shaders/fft.wgsl"), + &fft_bgl, + ); + eprintln!( + "[sonar_wgpu] [{label}] GPU pipelines compiled in {:.0} ms -> ready.", + t0.elapsed().as_millis() + ); + + Some(GpuContext { + device, + queue, + bs_pipeline, + bs_bgl, + convert_pipeline, + convert_bgl, + mm_pipeline, + mm_bgl, + fft_pipeline, + fft_bgl, + buffers: std::sync::Mutex::new(HashMap::new()), + }) + }); + + match result { + Ok(ctx) => ctx, + Err(_) => { + eprintln!("[sonar_wgpu] [{label}] GPU init panicked"); + None + } + } +} + +/// GPU context initialization: adapter enumeration, device request. +fn init_gpu_context() -> Option { + if let Some(backends) = wgpu::util::backend_bits_from_env() { + eprintln!("[sonar_wgpu] WGPU_BACKEND set; using only {backends:?}"); + return try_init_gpu_context(backends, "WGPU_BACKEND"); + } + for (backends, label) in BACKEND_ATTEMPTS { + if let Some(ctx) = try_init_gpu_context(*backends, label) { + return Some(ctx); + } + } + + eprintln!("[sonar_wgpu] GPU init failed for all backend sets"); + None +} + +/// Global singleton: lazily initializes GPU context on first access. +/// Returns None if no usable GPU adapter was found. +pub fn get_or_init() -> Option<&'static GpuContext> { + GPU_CONTEXT.get_or_init(init_gpu_context).as_ref() +} diff --git a/gazebo/posim_gz_multibeam_sonar/wgpu_vendor/sonar_wgpu_rust/src/shaders/backscatter.wgsl b/gazebo/posim_gz_multibeam_sonar/wgpu_vendor/sonar_wgpu_rust/src/shaders/backscatter.wgsl new file mode 100644 index 00000000..a3a7a643 --- /dev/null +++ b/gazebo/posim_gz_multibeam_sonar/wgpu_vendor/sonar_wgpu_rust/src/shaders/backscatter.wgsl @@ -0,0 +1,179 @@ +/// The physics kernel. + +struct Params { + n_beams: u32, + n_rays: u32, + n_freq: u32, + ray_skips: u32, + sound_speed: f32, + max_distance: f32, + source_level: f32, + attenuation: f32, + sensor_gain: f32, + bandwidth: f32, + seed_lo: u32, + seed_hi: u32, + frame_lo: u32, + frame_hi: u32, + area_scaler: f32, + h_fov: f32, + v_fov: f32, + _pad0: u32, + _pad1: u32, + _pad2: u32, +}; + +@group(0) @binding(0) var params: Params; +@group(0) @binding(1) var depth: array; +@group(0) @binding(2) var normals: array; +@group(0) @binding(3) var reflectivity: array; +@group(0) @binding(4) var out_re: array>; +@group(0) @binding(5) var out_im: array>; +@group(0) @binding(6) var window_vals: array; + +const PI: f32 = 3.14159265358979323846; +const SCALE: f32 = 1024.0; +const LIMIT: i32 = 2147483520i; // near i32 max + +// Philox4x32-10 RNG: WGPU/WGSL port of CUDA curand implementation. +// Based on Salmon et al., 2011, SC'11: "Random123: A Library of Counter-Based Random Number Generators". +// Based on https://github.com/DEShawResearch/random123/blob/main/include/Random123/philox.h +fn mulhilo32(a: u32, b: u32) -> vec2 { + let a_lo = a & 0xFFFFu; + let a_hi = a >> 16u; + let b_lo = b & 0xFFFFu; + let b_hi = b >> 16u; + let lo = a * b; + let p0 = a_lo * b_lo; + let p1 = a_hi * b_lo; + let p2 = a_lo * b_hi; + let p3 = a_hi * b_hi; + let mid = (p0 >> 16u) + (p1 & 0xFFFFu) + (p2 & 0xFFFFu); + let hi = p3 + (p1 >> 16u) + (p2 >> 16u) + (mid >> 16u); + return vec2(hi, lo); +} + +// Single Philox4x32 round. +fn philox_round(c: vec4, k: vec2) -> vec4 { + let r0 = mulhilo32(0xD2511F53u, c.x); + let r1 = mulhilo32(0xCD9E8D57u, c.z); + return vec4(r1.x ^ c.y ^ k.x, r1.y, r0.x ^ c.w ^ k.y, r0.y); +} + +// Full 10-round Philox4x32. +fn philox4x32_10(ctr: vec4, key: vec2) -> vec4 { + var c = ctr; + var k = key; + c = philox_round(c, k); k.x += 0x9E3779B9u; k.y += 0xBB67AE85u; + c = philox_round(c, k); k.x += 0x9E3779B9u; k.y += 0xBB67AE85u; + c = philox_round(c, k); k.x += 0x9E3779B9u; k.y += 0xBB67AE85u; + c = philox_round(c, k); k.x += 0x9E3779B9u; k.y += 0xBB67AE85u; + c = philox_round(c, k); k.x += 0x9E3779B9u; k.y += 0xBB67AE85u; + c = philox_round(c, k); k.x += 0x9E3779B9u; k.y += 0xBB67AE85u; + c = philox_round(c, k); k.x += 0x9E3779B9u; k.y += 0xBB67AE85u; + c = philox_round(c, k); k.x += 0x9E3779B9u; k.y += 0xBB67AE85u; + c = philox_round(c, k); k.x += 0x9E3779B9u; k.y += 0xBB67AE85u; + c = philox_round(c, k); + return c; +} + +// Convert raw u32 to uniform float (0,1]. +fn philox_uniform(x: u32) -> f32 { + return f32(x) * (1.0 / 4294967296.0) + (0.5 / 4294967296.0); +} + +// Generate 4 normal-distributed floats using Box-Muller transform. +// Counter: [frame_lo, frame_hi, subsequence, 0], Key: [seed_lo, seed_hi] +// Mirrors CUDA curand_init(seed, subsequence, offset=0) +fn philox_normal4(seed_lo: u32, seed_hi: u32, frame_lo: u32, frame_hi: u32, subsequence: u32) -> vec4 { + let raw = philox4x32_10(vec4(frame_lo, frame_hi, subsequence, 0u), vec2(seed_lo, seed_hi)); + let u1 = philox_uniform(raw.x); + let u2 = philox_uniform(raw.y); + let u3 = philox_uniform(raw.z); + let u4 = philox_uniform(raw.w); + let r1 = sqrt(-2.0 * log(u1)); + let t1 = 2.0 * PI * u2; + let r2 = sqrt(-2.0 * log(u3)); + let t2 = 2.0 * PI * u4; + return vec4(r1 * sin(t1), r1 * cos(t1), r2 * sin(t2), r2 * cos(t2)); +} + +@compute @workgroup_size(8, 8, 1) +fn main(@builtin(global_invocation_id) gid: vec3) { + let beam = gid.x; + let reduced_ray = gid.y; + if (beam >= params.n_beams) { + return; + } + + let ray_step = max(1u, params.ray_skips); + let actual_ray = reduced_ray * ray_step; + if (actual_ray >= params.n_rays) { + return; + } + + let dr = beam * params.n_rays + actual_ray; + let d = depth[dr]; + if (d <= 0.0 || d > params.max_distance) { + return; + } + + let nx = normals[dr * 3u + 0u]; + let ny = normals[dr * 3u + 1u]; + let nz = normals[dr * 3u + 2u]; + let refl = reflectivity[dr]; + + // Compute sensor-frame ray unit direction from (beam, actual_ray) pixel indices. + // Direction formula (sin(θ_b)cos(θ_r), sin(θ_r), cos(θ_b)cos(θ_r)) is unit-length by construction. + let beam_ang = params.h_fov * (f32(beam) / max(f32(params.n_beams) - 1.0, 1.0) - 0.5); + let ray_ang = params.v_fov * (f32(actual_ray) / max(f32(params.n_rays) - 1.0, 1.0) - 0.5); + let rd_x = sin(beam_ang) * cos(ray_ang); + let rd_y = sin(ray_ang); + let rd_z = cos(beam_ang) * cos(ray_ang); + // let cos_inc = abs(rd_x * nx + rd_y * ny + rd_z * nz); + let cos_inc = abs(nz); + let lambert_sqrt = sqrt(max(refl, 0.0)) * cos_inc; + + let source_term = sqrt(pow(10.0, params.source_level / 10.0)) * 1e-6; + let propagation = exp(-2.0 * params.attenuation * d) / (d * d); + let target_area_sqrt = sqrt(d * params.area_scaler); + // let amp_scale = source_term * propagation * lambert_sqrt * target_area_sqrt * params.sensor_gain; + let amp_scale = source_term * propagation * lambert_sqrt * target_area_sqrt; + + let subsequence = beam * params.n_rays + actual_ray; + let xi = philox_normal4(params.seed_lo, params.seed_hi, params.frame_lo, params.frame_hi, subsequence); + + let amp_re = (xi.x / sqrt(2.0)) * amp_scale; + let amp_im = (xi.y / sqrt(2.0)) * amp_scale; + + let delta_f = params.bandwidth / f32(params.n_freq); + let n_freq_f = f32(params.n_freq); + let is_even = (params.n_freq % 2u) == 0u; + + // Keep window buffer bound for interface stability. + _ = window_vals[0u]; + + // Per-frequency phase sweep + for (var f: u32 = 0u; f < params.n_freq; f = f + 1u) { + var freq: f32; + if (is_even) { + freq = delta_f * (-n_freq_f + 2.0 * (f32(f) + 1.0)) / 2.0; + } else { + freq = delta_f * (-(n_freq_f - 1.0) + 2.0 * (f32(f) + 1.0)) / 2.0; + } + + let kw = 2.0 * PI * freq / params.sound_speed; + let phase = 2.0 * d * kw; + let c = cos(phase); + let s = sin(phase); + + // let re_val = clamp((amp_re * c - amp_im * s) * SCALE, -LIMIT, LIMIT); + // let im_val = clamp((amp_re * s + amp_im * c) * SCALE, -LIMIT, LIMIT); + let re_val = clamp(i32(round((amp_re * c - amp_im * s) * SCALE)), -LIMIT, LIMIT); + let im_val = clamp(i32(round((amp_re * s + amp_im * c) * SCALE)), -LIMIT, LIMIT); + + let idx = beam * params.n_freq + f; + atomicAdd(&out_re[idx], i32(re_val)); + atomicAdd(&out_im[idx], i32(im_val)); + } +} diff --git a/gazebo/posim_gz_multibeam_sonar/wgpu_vendor/sonar_wgpu_rust/src/shaders/convert.wgsl b/gazebo/posim_gz_multibeam_sonar/wgpu_vendor/sonar_wgpu_rust/src/shaders/convert.wgsl new file mode 100644 index 00000000..84f66af6 --- /dev/null +++ b/gazebo/posim_gz_multibeam_sonar/wgpu_vendor/sonar_wgpu_rust/src/shaders/convert.wgsl @@ -0,0 +1,15 @@ +/// Fixed-point atomic accumulation in backscatter.wgsl (SCALE=1024) allows us +/// to avoid floating-point atomics, which are not widely supported on GPUs. This shader +/// converts the accumulated i32 values back to f32 by dividing by SCALE=1024, precision 0.001. +@group(0) @binding(0) var in_re: array; +@group(0) @binding(1) var in_im: array; +@group(0) @binding(2) var out_re: array; +@group(0) @binding(3) var out_im: array; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { + let i = gid.x; + if (i >= arrayLength(&in_re)) { return; } + out_re[i] = f32(in_re[i]) / 1024.0; + out_im[i] = f32(in_im[i]) / 1024.0; +} diff --git a/gazebo/posim_gz_multibeam_sonar/wgpu_vendor/sonar_wgpu_rust/src/shaders/fft.wgsl b/gazebo/posim_gz_multibeam_sonar/wgpu_vendor/sonar_wgpu_rust/src/shaders/fft.wgsl new file mode 100644 index 00000000..d56e6b23 --- /dev/null +++ b/gazebo/posim_gz_multibeam_sonar/wgpu_vendor/sonar_wgpu_rust/src/shaders/fft.wgsl @@ -0,0 +1,107 @@ +// Cooley-Tukey FFT with zero-padding to the next power of 2. +// Supports any n_freq <= 4096. The input is padded to padded_n (next power of 2 >= n_freq) +// in shared memory. The padded FFT is then sampled back onto the original n_freq +// physical range grid before writeback. + +struct Params { + n_beams: u32, + n_freq: u32, // actual number of frequency bins (may not be power-of-2) + padded_n: u32, // next power of 2 >= n_freq (actual FFT size, <= 4096) + log2_n: u32, // log2(padded_n) +}; + +@group(0) @binding(0) var params: Params; +@group(0) @binding(1) var p_re: array; +@group(0) @binding(2) var p_im: array; + +// Shared scratchpad for in-place FFT: sized for the maximum supported padded_n. +var smem_re: array; +var smem_im: array; + +// Bit-reversal permutation used by the Cooley-Tukey butterfly. +fn bit_reverse(v: u32, bits: u32) -> u32 { + var x = v; + var y: u32 = 0u; + for (var i: u32 = 0u; i < bits; i = i + 1u) { + y = (y << 1u) | (x & 1u); + x = x >> 1u; + } + return y; +} + +// One workgroup per beam. Threads: 256 per workgroup. +@compute @workgroup_size(256, 1, 1) +fn main( + @builtin(workgroup_id) wg: vec3, // wg.x = beam index + @builtin(local_invocation_id) lid: vec3 // lid.x = thread 0..255 +) { + let beam = wg.x; + if (beam >= params.n_beams || params.padded_n > 4096u) { + return; + } + + let n = params.padded_n; // FFT size (power of 2) + let n_real = params.n_freq; // actual data count stored in the buffer + let base = beam * n_real; // flat offset into the (non-padded) p_re/p_im buffers + + // Load with zero-padding into bit-reversed positions: + // indices [0, n_real) come from the buffer; [n_real, n) are zero-padded. + for (var i: u32 = lid.x; i < n; i = i + 256u) { + let j = bit_reverse(i, params.log2_n); + if (i < n_real) { + smem_re[j] = p_re[base + i]; + smem_im[j] = p_im[base + i]; + } else { + smem_re[j] = 0.0; + smem_im[j] = 0.0; + } + } + workgroupBarrier(); + + // Cooley-Tukey butterfly iterations over log2(padded_n) stages. + for (var s: u32 = 0u; s < params.log2_n; s = s + 1u) { + let half = 1u << s; + let stride = half << 1u; + + for (var i: u32 = lid.x; i < n / 2u; i = i + 256u) { + let group = i / half; + let j = i % half; + let i0 = group * stride + j; + let i1 = i0 + half; + + // Twiddle factor: exp(-j * 2*pi * k / stride) + let angle = -2.0 * 3.14159265358979323846 * f32(j) / f32(stride); + let wr = cos(angle); + let wi = sin(angle); + + // Butterfly: u ± w*t + let tr = smem_re[i1] * wr - smem_im[i1] * wi; + let ti = smem_re[i1] * wi + smem_im[i1] * wr; + let ur = smem_re[i0]; + let ui = smem_im[i0]; + + smem_re[i0] = ur + tr; + smem_im[i0] = ui + ti; + smem_re[i1] = ur - tr; + smem_im[i1] = ui - ti; + } + + workgroupBarrier(); + } + + // Range-preserving writeback. + // + // A target at original-grid bin k appears at approximately + // k * padded_n / n_real in the padded FFT. Copying smem[k] directly therefore + // stretches the reported range by padded_n / n_real and drops the far-range + // tail. Sample the padded spectrum at that scaled coordinate and linearly + // interpolate the complex components back to n_real output bins. + for (var i: u32 = lid.x; i < n_real; i = i + 256u) { + let src = f32(i) * f32(n) / f32(n_real); + let lo = min(u32(floor(src)), n - 1u); + let hi = min(lo + 1u, n - 1u); + let alpha = src - f32(lo); + p_re[base + i] = mix(smem_re[lo], smem_re[hi], alpha); + p_im[base + i] = mix(smem_im[lo], smem_im[hi], alpha); + } +} diff --git a/gazebo/posim_gz_multibeam_sonar/wgpu_vendor/sonar_wgpu_rust/src/shaders/matmul.wgsl b/gazebo/posim_gz_multibeam_sonar/wgpu_vendor/sonar_wgpu_rust/src/shaders/matmul.wgsl new file mode 100644 index 00000000..b1e27318 --- /dev/null +++ b/gazebo/posim_gz_multibeam_sonar/wgpu_vendor/sonar_wgpu_rust/src/shaders/matmul.wgsl @@ -0,0 +1,71 @@ +// GPU GEMM: WGPU compute shader for beam correction via matrix multiply. +// Tiled approach: C = A*B normalized, with shared memory for cache optimization. + +struct Params { + n_beams: u32, + n_freq: u32, + beam_corrector_sum: f32, +}; + +@group(0) @binding(0) var params: Params; +@group(0) @binding(1) var a: array; // beam_corrector [nBeams*nBeams] +@group(0) @binding(2) var b: array; // spectrum input [nBeams*nFreq] +@group(0) @binding(3) var c: array; // output [nBeams*nFreq] + +// Shared memory tiles for GEMM optimization (tile size 16x16). +var tile_a: array, 16>; +var tile_b: array, 16>; + +// Tiled GEMM kernel: 16x16 workgroup, each thread computes one C[row,col] element. +@compute @workgroup_size(16, 16, 1) +fn main( + @builtin(global_invocation_id) gid: vec3, + @builtin(local_invocation_id) lid: vec3 +) { + let row = gid.y; + let col = gid.x; + let in_bounds = (row < params.n_beams) && (col < params.n_freq); + + // Accumulator: accumulates partial products across tiles. + var acc = 0.0; + let tiles = (params.n_beams + 15u) / 16u; + + // Loop over tiles of size 16: classic blocked GEMM. + // (k = beam index in the inner product) + for (var t: u32 = 0u; t < tiles; t = t + 1u) { + let k_a = t * 16u + lid.x; + let k_b = t * 16u + lid.y; + + // LOAD: load tile data into shared memory. + // Load A tile: each thread loads one A[row, k_a]. + if (in_bounds && k_a < params.n_beams) { + tile_a[lid.y][lid.x] = a[row * params.n_beams + k_a]; + } else { + tile_a[lid.y][lid.x] = 0.0; + } + + // Load B tile: each thread loads one B[k_b, col]. + if (in_bounds && k_b < params.n_beams) { + tile_b[lid.y][lid.x] = b[k_b * params.n_freq + col]; + } else { + tile_b[lid.y][lid.x] = 0.0; + } + + workgroupBarrier(); + + // COMPUTE: partial dot product. + // Accumulate tile product: inner loop over tile dimension. + for (var k: u32 = 0u; k < 16u; k = k + 1u) { + acc = acc + tile_a[lid.y][k] * tile_b[k][lid.x]; + } + + workgroupBarrier(); + } + + // WRITE: final result to global memory. + // Normalize output: divide by beam_corrector_sum. + let norm = max(params.beam_corrector_sum, 1e-12); + if (in_bounds) { + c[row * params.n_freq + col] = acc / norm; + } +} diff --git a/models/posim_sensor_models/description/blueview_p900/model.sdf b/models/posim_sensor_models/description/blueview_p900/model.sdf index dff61f43..01d96527 100644 --- a/models/posim_sensor_models/description/blueview_p900/model.sdf +++ b/models/posim_sensor_models/description/blueview_p900/model.sdf @@ -85,6 +85,8 @@ 12 900e3 29.9e3 + + 1500 220 10 diff --git a/website/catalog.json b/website/catalog.json index 34924ddc..c3ec48d0 100644 --- a/website/catalog.json +++ b/website/catalog.json @@ -1,5 +1,5 @@ { - "sourceRevision": "23bb608a9a89b9cda3472792316a81711f1ce73e", + "sourceRevision": "ad6165bd6c2b4e407f18ecb264e20bb76da8d4ca", "worlds": [ "camera_tutorial.world", "dvl_world.world", diff --git a/website/content/en/index.md b/website/content/en/index.md index d5dad934..ed432244 100644 --- a/website/content/en/index.md +++ b/website/content/en/index.md @@ -29,7 +29,7 @@ Prepare an [Ubuntu installation](install.md) or [Docker environment](docker.md), | --- | --- | --- | | Ocean worlds | Surface, seabed and task scenes | [World library](worlds.md) | | Robots | REXROV, BlueROV2, Slocum and surface vehicles | [ROVs](rovs.md), [Gliders](gliders.md), [Surface robots](surface.md) | -| Sensors | Camera, DVL, pressure, USBL and CUDA sonar | [Camera](camera.md), [DVL](dvl.md), [Sonar](sonar.md) | +| Sensors | Camera, DVL, pressure, USBL and multibeam sonar | [Camera](camera.md), [DVL](dvl.md), [Sonar](sonar.md) | | Environment plugins | Currents and geographic coordinates | [Currents](currents.md), [Coordinates](coordinates.md) | | ROS interfaces | Sensor subscriptions, commands and recording | [ROS 2 and control](ros.md) | diff --git a/website/content/en/install.md b/website/content/en/install.md index fb5681e6..e2880e72 100644 --- a/website/content/en/install.md +++ b/website/content/en/install.md @@ -47,6 +47,9 @@ rosdep install --rosdistro lyrical --from-paths src --ignore-src -r -y ## 4. Build and source ```bash +colcon build --merge-install --executor sequential --symlink-install \ + --packages-select wgpu_vendor +source install/setup.bash colcon build --merge-install --executor sequential --symlink-install source install/setup.bash ros2 pkg prefix posim_demos @@ -60,7 +63,7 @@ source "$HOME/.ros_ardusub_env/env" source ~/posim_ws/install/setup.bash ``` -Continue with [First simulation](quickstart.md). NVIDIA/CUDA sonar has additional [build requirements](sonar-tuning.md). +Continue with [First simulation](quickstart.md). Choose the sonar compute backend using [Sonar build and performance](sonar-tuning.md). If the compiler is killed because memory is exhausted, limit compiler jobs as well as package concurrency. `--executor sequential` serializes packages, not compiler processes: diff --git a/website/content/en/requirements.md b/website/content/en/requirements.md index 30f7a4df..9b66d7fe 100644 --- a/website/content/en/requirements.md +++ b/website/content/en/requirements.md @@ -21,7 +21,7 @@ Run the native recipe in a Bash shell on Ubuntu. Docker commands run on the host Camera, depth and DVL processing require a renderer even without a desktop window. Software rendering can provide an initial data check. -Multibeam sonar requires NVIDIA GPU access, CUDA and built sonar libraries. See [Sonar build and performance](sonar-tuning.md). +Multibeam sonar supports WGPU, CUDA and CPU processing. Build its Rust/Cargo library and sonar plugins, then prepare the selected backend and Gazebo rendering. See [Sonar build and performance](sonar-tuning.md). ## Plan resources for your experiment diff --git a/website/content/en/sonar-tuning.md b/website/content/en/sonar-tuning.md index e7a4b6c8..f7f08738 100644 --- a/website/content/en/sonar-tuning.md +++ b/website/content/en/sonar-tuning.md @@ -1,46 +1,121 @@ # Sonar build and performance -Multibeam sonar uses CUDA. Prepare the NVIDIA driver, CUDA compiler/runtime, cuFFT, cuBLAS and Gazebo rendering. +Build the multibeam sonar once, then select WGPU, CUDA or CPU when launching it. WGPU uses Vulkan on Linux and Metal on native macOS. All compute choices require Gazebo rendering for the ray sensor. -## Check the toolchain +## Prepare the toolchain -On the machine or inside the container that will run sonar: +Use the [POSIM workspace setup](install.md) with ROS 2 Lyrical, Gazebo Jetty and ROS–Gazebo bridge packages. Run the build from the workspace root, such as `~/posim_ws`. + +- Rust and Cargo are required to build the WGPU library, including when you plan to select CPU processing. +- Linux WGPU needs a Vulkan loader and an appropriate GPU driver. `vulkaninfo --summary` lists available adapters. +- CUDA needs a compatible NVIDIA GPU, its driver, the CUDA compiler/runtime, cuFFT and cuBLAS. CUDA support is included when CMake finds the toolkit. +- Rendering needs Ogre2 and a working display, virtual display or configured headless renderer. + +```bash +rustc --version +cargo --version +``` + +If Rust is unavailable, use the [official Rust installer](https://www.rust-lang.org/tools/install). For CUDA, also check: ```bash nvidia-smi nvcc --version ``` -GPU visibility and compiler availability are separate checks. For containers, configure [NVIDIA Container Toolkit](https://docs.nvidia.com/datacenter/cloud-native/container-toolkit/latest/install-guide.html) on the host as well as the required libraries inside the image. +GPU access and compiler availability are separate checks. Linux containers using NVIDIA GPUs also need the host's [NVIDIA Container Toolkit](https://docs.nvidia.com/datacenter/cloud-native/container-toolkit/latest/install-guide.html) and the driver/runtime libraries inside the container. + +## Build and source + +On a fresh Ubuntu workspace, build and source the WGPU library before building the remaining packages: + +```bash +source /opt/ros/lyrical/setup.bash +colcon build --merge-install --executor sequential --symlink-install \ + --packages-select wgpu_vendor +source install/setup.bash +colcon build --merge-install --executor sequential --symlink-install +source install/setup.bash +``` + +In an already sourced native macOS workspace, use the same two build commands with that workspace's ROS environment. Rebuild the workspace after changing its compiler or GPU toolchain. ## Select a CUDA architecture -The sonar CMake configuration exposes `CUDA_ARCHITECTURE`, defaulting to `60`. Set it to the SM version supported by your GPU and CUDA toolkit. From a sourced workspace, rebuild after choosing that value: +The sonar CMake variable `CUDA_ARCHITECTURE` defaults to `native`. It selects the local GPU's architecture when CUDA support is built. For a different deployment GPU, supply the SM value supported by that GPU and toolkit: ```bash colcon build --merge-install --executor sequential --symlink-install \ + --packages-select multibeam_sonar \ --cmake-args -DCUDA_ARCHITECTURE=YOUR_GPU_SM source install/setup.bash ``` -Replace `YOUR_GPU_SM` with the numeric architecture. Check CUDA's supported architectures before selecting it; a toolkit may have dropped an older target. Do not confuse this project-specific singular variable with CMake's `CMAKE_CUDA_ARCHITECTURES`. +Replace `YOUR_GPU_SM` with the numeric architecture. The project forwards this value to CMake's `CMAKE_CUDA_ARCHITECTURES`. Check the toolkit's supported targets before selecting a value. ## Verify installed targets ```bash +ros2 pkg prefix wgpu_vendor ros2 pkg prefix multibeam_sonar ros2 pkg prefix multibeam_sonar_system ros2 pkg prefix posim_multibeam_sonar_demo ``` -Check library files to confirm CUDA targets were built. Packages can register even when CUDA targets are skipped. +On Linux, check the installed libraries: ```bash +test -f "$(ros2 pkg prefix wgpu_vendor)/lib/libsonar_wgpu.a" test -f "$(ros2 pkg prefix multibeam_sonar)/lib/multibeam_sonar/libmultibeam_sonar.so" test -f "$(ros2 pkg prefix multibeam_sonar_system)/lib/multibeam_sonar_system/libmultibeam_sonar_system.so" ``` -Each command must exit with status 0. Check build/launch logs for skipped targets or load errors, then receive image/raw data in [Multibeam sonar](sonar.md). +Each command must exit with status 0. Run the [sonar example](sonar.md) and confirm the selected backend in its log, then receive image and raw data. A package prefix or library file alone does not confirm GPU access. + +## Choose a WGPU adapter + +Set these variables in the terminal that launches Gazebo: + +| Variable | Behavior | +| --- | --- | +| `WGPU_BACKEND` | Selects the GPU API, such as `vulkan` on Linux or `metal` on native macOS | +| `WGPU_ADAPTER_NAME` | Selects the first adapter whose name contains the value, ignoring case | +| `WGPU_POWER_PREF` | Requests `high`, `low` or `none`; the default is `high` | + +For example, select a Vulkan adapter whose name contains `RTX`: + +```bash +WGPU_BACKEND=vulkan WGPU_ADAPTER_NAME=RTX WGPU_POWER_PREF=high \ +ros2 launch posim_multibeam_sonar_demo multibeam_sonar_demo.launch.py \ + compute_backend:=wgpu +``` + +Use an adapter name available on your machine. The log lists available names when a requested name has no match. Once selected, the WGPU adapter is shared by the sonar sensors in that Gazebo process; start a new process to change it. + +## Apple Silicon + +For native macOS execution, prepare a sourced ROS 2 workspace with Gazebo Jetty, Ogre2 and the ROS–Gazebo bridge. Build Rust/Cargo and the sonar packages as above. The WGPU compute backend uses the Apple GPU through Metal: + +```bash +WGPU_BACKEND=metal \ +ros2 launch posim_multibeam_sonar_demo multibeam_sonar_demo.launch.py \ + compute_backend:=wgpu +``` + +Metal provides the acoustic computation; Gazebo still needs its rendering environment. Use `compute_backend:=cpu` to select CPU acoustic processing. CUDA requires a compatible NVIDIA environment. + +For a Linux ARM64 container on an Apple Silicon host, follow [Docker setup and virtual-display rendering](docker.md). Run the [server-only sonar command](sonar.md) with `POSIM_SONAR_COMPUTE_BACKEND=cpu` inside the container. A Linux container requires a Vulkan adapter for WGPU; its compute API is separate from native macOS Metal. + +## Resolve startup errors + +| Symptom | Action | +| --- | --- | +| `wgpu_vendorConfig.cmake` is missing | Build `wgpu_vendor`, source `install/setup.bash`, then build the remaining packages | +| WGPU reports no usable adapter | Check the GPU driver, `vulkaninfo --summary` on Linux and adapter-selection variables; choose `auto` or `cpu` if GPU processing is unavailable | +| Explicit CUDA selection fails | Check CUDA build support, `nvidia-smi`, the runtime libraries and container GPU access | +| An unknown backend is reported | Use `auto`, `wgpu`, `cuda` or `cpu` | +| Another world with the same name is running | Stop the existing POSIM launch with Ctrl+C before relaunching | +| Backend initializes but no image arrives | Confirm the world is unpaused, then check the renderer and the [sonar topics](sonar.md) | ## Tune one variable at a time @@ -53,12 +128,12 @@ Each command must exit with status 0. Check build/launch logs for skipped target | `writeLog`, `debugFlag` | Disk and console overhead | | Sensor gain | Image presentation; distinguish it from physical return strength | -Record frame processing time, received topic rate, Gazebo real-time factor, CPU/GPU utilization and memory with the same scene and pose. Separate physics, rendering, CUDA computation, publication and logging costs. Warm up the renderer before timing. +Record frame processing time, received topic rate, Gazebo real-time factor and CPU/GPU utilization in the same scene and pose. Separate physics, rendering, acoustic computation, publication and recording costs. Allow the renderer to warm up before measuring. The configured 30 Hz update rate is a request; received throughput depends on the scene, sensor settings and hardware. -Retain a baseline dataset when changing numerical kernels. Compare intensity/range outputs, not only rendered screenshots. Upstream research and attribution are listed under [Citation](citation.md). +Use raw intensity–range data when comparing backends or sensor settings, and keep gain and display settings fixed. Research sources and attribution are listed under [Citation](citation.md). ## GPU processing flow
GPU stages in sonar processing
GPU stages in sonar processing
-Sonar processing includes ray summation, window/correction and FFT stages. Separate stage timings from GPU memory transfers when profiling. Figure: POSIM Notion Wiki. +Sonar processing includes backscatter, beam correction and range compression. Separate computation time from memory transfers when profiling. Figure: POSIM Notion Wiki. diff --git a/website/content/en/sonar.md b/website/content/en/sonar.md index 93c7ac57..5550c790 100644 --- a/website/content/en/sonar.md +++ b/website/content/en/sonar.md @@ -1,22 +1,41 @@ # Multibeam sonar -POSIM's multibeam sonar uses a ray-based point-scattering model to generate intensity–range data and sonar imagery. Prepare the [NVIDIA/CUDA backend](sonar-tuning.md) before running the sensor. +POSIM's multibeam sonar generates intensity–range data, sonar images and point clouds from a ray-based point-scattering model. Choose WGPU, CUDA or CPU processing at launch. Prepare the [sonar build and rendering environment](sonar-tuning.md) first. ## Run a sonar example +From a sourced POSIM workspace, launch the BlueView P900 demo with WGPU: + ```bash -ros2 launch posim_demos posim_sensor.launch.py \ - namespace:=blueview_p900 world_name:=posim_multibeam_sonar \ - paused:=false x:=4 z:=2.0 yaw:=3.14 gui:=false headless:=true +ros2 launch posim_multibeam_sonar_demo multibeam_sonar_demo.launch.py \ + compute_backend:=wgpu ``` -A dedicated example is also supplied by `posim_multibeam_sonar_demo`: +The demo opens Gazebo and RViz. RViz shows the sensor's point cloud. Use `rviz:=false` to run the demo without the RViz window. Run one launch at a time and stop it with Ctrl+C before starting another. + +## Select a compute backend + +Change `compute_backend` when launching the demo. The installed plugin supports these choices without recompiling: + +| Value | Selection and prerequisites | +| --- | --- | +| `auto` (default) | Tries CUDA, then WGPU, then CPU | +| `wgpu` | Uses a WGPU adapter: Vulkan on Linux, Metal on native macOS | +| `cuda` | Uses an NVIDIA GPU; requires CUDA support in the plugin and a working driver/runtime | +| `cpu` | Computes acoustic returns on the CPU; Gazebo rendering is still required | + +Check the launch log for `Using sonar compute backend:` and the selected WGPU adapter. Explicit `wgpu` or `cuda` selection requires that backend to initialize. An unavailable backend or an invalid name disables that sonar sensor and reports an error. Choose `auto` to allow backend fallback, or `cpu` to request CPU processing directly. + +For the general sensor launch, set `POSIM_SONAR_COMPUTE_BACKEND` in the launching environment: ```bash -ros2 launch posim_multibeam_sonar_demo multibeam_sonar_demo.launch.py +POSIM_SONAR_COMPUTE_BACKEND=cpu \ +ros2 launch posim_demos posim_sensor.launch.py \ + namespace:=blueview_p900 world_name:=posim_multibeam_sonar \ + paused:=false x:=4 z:=2.0 yaw:=3.14 gui:=false headless:=true ``` -Use one launch at a time. For the robot variant, select `bluerov2_heavy_multibeam_sonar` with a suitable sonar world. +Server mode also needs a rendering environment. Follow the [Docker rendering instructions](docker.md) for a virtual display or configured GPU access. For the robot variant, select `bluerov2_heavy_multibeam_sonar` with a suitable sonar world. ## Observe outputs @@ -24,9 +43,20 @@ Use one launch at a time. For the robot variant, select `bluerov2_heavy_multibea ros2 topic list -t ros2 topic echo /sensor/multibeam_sonar/sonar_image \ sensor_msgs/msg/Image --once --no-arr +ros2 topic hz /sensor/multibeam_sonar/sonar_image ``` -In RViz, add an Image display for the sonar image and PointCloud2 for `/sensor/multibeam_sonar/point_cloud`. Raw returns use the acoustic message interface supplied by `marine_acoustic_msgs`; discover their type with `ros2 topic list -t`. +| Topic | ROS 2 message | Use | +| --- | --- | --- | +| `/sensor/multibeam_sonar/sonar_image` | `sensor_msgs/msg/Image` | Display sonar intensity in RViz's Image display | +| `/sensor/multibeam_sonar/sonar_image_raw` | `marine_acoustic_msgs/msg/ProjectedSonarImage` | Read acoustic observations with range and beam information | +| `/sensor/multibeam_sonar/point_cloud` | `sensor_msgs/msg/PointCloud2` | Display the ray geometry in RViz's PointCloud2 display | + +The sonar publishes image and raw acoustic messages directly to ROS 2. The sensor configuration bridges the Gazebo point cloud to ROS 2. Discover custom topic names with `ros2 topic list -t`. + +## How the sensor works + +Gazebo supplies ray depths and surface normals for the scene. The sonar model computes acoustic backscatter, applies beam correction and compresses the returns into range bins. The resulting acoustic data and display image describe the same observation. Use raw data for analysis and retain the gain and display settings with recorded images. ## Sensor configuration @@ -51,7 +81,7 @@ The `` block controls `sonarFreq`, `bandwidth`, `soundSpeed`, `sourceLevel ## Add a sonar sensor and world -Copy `models/posim_sensor_models/description/blueview_p900/` to a new sensor identifier. Also prepare a matching `config//sensor_config.py`; align SDF model/link/frame names with bridge topics. Set horizontal/vertical samples, angles and range in `` to match your sensor specification. +Copy `models/posim_sensor_models/description/blueview_p900/` to a new sensor identifier. Also prepare a matching `config//sensor_config.py`; align SDF model/link/frame names with bridge topics. Set horizontal beams, vertical rays, angles and range in `` to match your sensor specification. Include the Gazebo rendering-sensor system and sonar system in the world. Consult the [sonar world SDF](https://github.com/IOES-Lab/POSIM/blob/main/models/posim_worlds/worlds/posim_multibeam_sonar.world) for physics, scene and other custom-sensor configuration. @@ -62,7 +92,7 @@ Include the Gazebo rendering-sensor system and sonar system in the world. Consul ``` -Prepare CUDA and rendering, rebuild/source the workspace, then select your descriptor using the [sensor launch arguments](quickstart.md). The ROS point-cloud bridge connects `sensor_msgs/msg/PointCloud2` to Gazebo `gz.msgs.PointCloudPacked`. Match `pointCloudTopicName` to the path configured in `sensor_config.py`. +Prepare the selected compute backend and rendering, rebuild/source the workspace, then select your descriptor using the [sensor launch arguments](quickstart.md). The ROS point-cloud bridge connects `sensor_msgs/msg/PointCloud2` to Gazebo `gz.msgs.PointCloudPacked`. Match `pointCloudTopicName` to the path configured in `sensor_config.py`. ## Record raw sonar data diff --git a/website/content/en/troubleshooting.md b/website/content/en/troubleshooting.md index b9c4c4f2..570e6109 100644 --- a/website/content/en/troubleshooting.md +++ b/website/content/en/troubleshooting.md @@ -17,7 +17,7 @@ Read the first missing URI in the log. Check that its package resources were ins ## No image or rendering initialization failure -Image sensors need a renderer even without a desktop window. Check display/offscreen settings, [Docker Xvfb setup](docker.md) and [CUDA libraries](sonar-tuning.md). +Image sensors need a renderer even without a desktop window. Check display/offscreen settings, [Docker Xvfb setup](docker.md) and [sonar backend setup](sonar-tuning.md). ## No messages diff --git a/website/content/ko/index.md b/website/content/ko/index.md index b4357428..c502d3af 100644 --- a/website/content/ko/index.md +++ b/website/content/ko/index.md @@ -29,7 +29,7 @@ POSIM은 [Project DAVE](https://field-robotics-lab.github.io/dave.doc/)를 계 | --- | --- | --- | | 해양 월드 | 수면·해저·작업 장면 | [월드 목록](worlds.md) | | 로봇 | REXROV, BlueROV2, Slocum, 수상 로봇 | [ROV](rovs.md), [글라이더](gliders.md), [수상 로봇](surface.md) | -| 센서 | 카메라·DVL·수압·USBL·CUDA 소나 | [카메라](camera.md), [DVL](dvl.md), [소나](sonar.md) | +| 센서 | 카메라·DVL·수압·USBL·멀티빔 소나 | [카메라](camera.md), [DVL](dvl.md), [소나](sonar.md) | | 환경 플러그인 | 해류·위도·경도 좌표 | [해류](currents.md), [좌표](coordinates.md) | | ROS 인터페이스 | 센서 구독·제어·실험 기록 | [ROS 2와 제어](ros.md) | diff --git a/website/content/ko/install.md b/website/content/ko/install.md index a49be70e..5f1c7080 100644 --- a/website/content/ko/install.md +++ b/website/content/ko/install.md @@ -47,6 +47,9 @@ rosdep install --rosdistro lyrical --from-paths src --ignore-src -r -y ## 4. 빌드하고 환경 불러오기 ```bash +colcon build --merge-install --executor sequential --symlink-install \ + --packages-select wgpu_vendor +source install/setup.bash colcon build --merge-install --executor sequential --symlink-install source install/setup.bash ros2 pkg prefix posim_demos @@ -60,7 +63,7 @@ source "$HOME/.ros_ardusub_env/env" source ~/posim_ws/install/setup.bash ``` -이어서 [첫 시뮬레이션](quickstart.md)을 실행하세요. NVIDIA/CUDA 소나는 [추가 빌드 설정](sonar-tuning.md)이 필요합니다. +이어서 [첫 시뮬레이션](quickstart.md)을 실행하세요. 소나 계산 방식은 [소나 빌드와 성능](sonar-tuning.md)에 따라 선택합니다. 메모리 부족으로 컴파일러가 강제 종료되면 패키지뿐 아니라 컴파일 작업 수도 줄이세요. `--executor sequential`은 패키지를 순서대로 빌드하지만, 패키지 안의 컴파일 작업 수까지 제한하지는 않습니다. diff --git a/website/content/ko/requirements.md b/website/content/ko/requirements.md index e9a80409..c519a7d1 100644 --- a/website/content/ko/requirements.md +++ b/website/content/ko/requirements.md @@ -21,7 +21,7 @@ Ubuntu 직접 설치 명령은 Bash에서 실행합니다. Docker 명령은 호 카메라·깊이 센서·DVL에는 렌더링 환경이 필요합니다. 화면 없이 실행할 때도 렌더러를 준비하세요. 첫 데이터 확인에는 소프트웨어 렌더링을 사용할 수 있습니다. -멀티빔 소나에는 NVIDIA GPU·CUDA 툴킷·소나 라이브러리가 필요합니다. [소나 빌드와 성능](sonar-tuning.md)을 참고하세요. +멀티빔 소나는 WGPU·CUDA·CPU 계산을 지원합니다. Rust·Cargo 라이브러리와 소나 플러그인을 빌드하고, 선택한 계산 방식과 Gazebo 렌더링을 준비하세요. [소나 빌드와 성능](sonar-tuning.md)을 참고하세요. ## 실험에 맞는 자원 준비 diff --git a/website/content/ko/sonar-tuning.md b/website/content/ko/sonar-tuning.md index 0ae58e27..1b5dd7e8 100644 --- a/website/content/ko/sonar-tuning.md +++ b/website/content/ko/sonar-tuning.md @@ -1,64 +1,139 @@ # 소나 빌드와 성능 -멀티빔 소나는 CUDA를 사용합니다. NVIDIA 드라이버·CUDA 컴파일러와 런타임·cuFFT·cuBLAS·Gazebo 렌더링을 준비하세요. +멀티빔 소나를 빌드한 뒤 실행할 때 WGPU·CUDA·CPU 계산 방식을 선택합니다. WGPU는 Linux에서 Vulkan, macOS에서 직접 실행할 때 Metal을 사용합니다. 모든 계산 방식에서 광선 센서의 Gazebo 렌더링이 필요합니다. -## 도구 확인 +## 도구 준비 -소나를 실행할 컴퓨터 또는 컨테이너 안에서 확인합니다. +[POSIM 작업 공간 설치 안내](install.md)에 따라 ROS 2 Lyrical·Gazebo Jetty·ROS–Gazebo 브리지를 준비합니다. `~/posim_ws`와 같은 작업 공간의 최상위 폴더에서 빌드하세요. + +- CPU 계산을 선택할 경우에도 WGPU 라이브러리를 빌드하기 위해 Rust와 Cargo가 필요합니다. +- Linux의 WGPU에는 Vulkan 로더와 GPU에 맞는 드라이버가 필요합니다. `vulkaninfo --summary`로 사용 가능한 어댑터를 확인합니다. +- CUDA에는 호환 NVIDIA GPU·드라이버·CUDA 컴파일러와 런타임·cuFFT·cuBLAS가 필요합니다. CMake가 툴킷을 찾으면 CUDA 지원을 포함합니다. +- 렌더링에는 Ogre2와 디스플레이·가상 디스플레이 또는 구성된 headless 렌더러가 필요합니다. + +```bash +rustc --version +cargo --version +``` + +Rust가 없으면 [공식 Rust 설치 안내](https://www.rust-lang.org/tools/install)를 따르세요. CUDA를 사용할 때는 다음 항목도 확인합니다. ```bash nvidia-smi nvcc --version ``` -GPU 접근과 컴파일러 설치는 각각 확인합니다. 컨테이너에서는 호스트의 [NVIDIA Container Toolkit](https://docs.nvidia.com/datacenter/cloud-native/container-toolkit/latest/install-guide.html)과 이미지 안의 라이브러리를 모두 준비해야 합니다. +GPU 접근과 컴파일러 설치는 각각 확인합니다. NVIDIA GPU를 사용하는 Linux 컨테이너에는 호스트의 [NVIDIA Container Toolkit](https://docs.nvidia.com/datacenter/cloud-native/container-toolkit/latest/install-guide.html)과 컨테이너 내부의 드라이버·런타임 라이브러리도 필요합니다. + +## 빌드하고 환경 불러오기 + +새 Ubuntu 작업 공간에서는 WGPU 라이브러리를 먼저 빌드하고 환경을 불러온 뒤 나머지 패키지를 빌드합니다. + +```bash +source /opt/ros/lyrical/setup.bash +colcon build --merge-install --executor sequential --symlink-install \ + --packages-select wgpu_vendor +source install/setup.bash +colcon build --merge-install --executor sequential --symlink-install +source install/setup.bash +``` + +macOS에서 직접 빌드할 때도 해당 작업 공간의 ROS 환경을 불러온 뒤 같은 두 빌드 명령을 사용합니다. 컴파일러나 GPU 도구를 바꾸면 작업 공간을 다시 빌드하세요. ## CUDA 아키텍처 선택 -소나 CMake에는 기본값이 `60`인 `CUDA_ARCHITECTURE` 변수가 있습니다. GPU와 CUDA 툴킷에서 지원하는 SM 버전을 선택한 뒤 작업 공간에서 다시 빌드합니다. +소나의 CMake 변수 `CUDA_ARCHITECTURE` 기본값은 `native`입니다. CUDA 지원을 빌드할 때 현재 GPU의 아키텍처를 선택합니다. 다른 GPU에 배포할 경우 해당 GPU와 툴킷이 지원하는 SM 값을 지정하세요. ```bash colcon build --merge-install --executor sequential --symlink-install \ + --packages-select multibeam_sonar \ --cmake-args -DCUDA_ARCHITECTURE=YOUR_GPU_SM source install/setup.bash ``` -`YOUR_GPU_SM`을 해당 아키텍처 숫자로 바꿉니다. 툴킷이 오래된 타깃 지원을 종료했을 수 있으므로 지원 목록을 확인하세요. 이 프로젝트의 단수형 변수와 CMake의 `CMAKE_CUDA_ARCHITECTURES`는 구분합니다. +`YOUR_GPU_SM`을 아키텍처 숫자로 바꿉니다. 프로젝트는 이 값을 CMake의 `CMAKE_CUDA_ARCHITECTURES`로 전달합니다. 값을 선택하기 전에 툴킷의 지원 목록을 확인하세요. ## 설치 타깃 확인 ```bash +ros2 pkg prefix wgpu_vendor ros2 pkg prefix multibeam_sonar ros2 pkg prefix multibeam_sonar_system ros2 pkg prefix posim_multibeam_sonar_demo ``` -CUDA 타깃이 빌드되었는지 라이브러리 파일을 확인합니다. 패키지는 CUDA 타깃을 생략해도 등록될 수 있습니다. +Linux에서는 설치된 라이브러리를 확인합니다. ```bash +test -f "$(ros2 pkg prefix wgpu_vendor)/lib/libsonar_wgpu.a" test -f "$(ros2 pkg prefix multibeam_sonar)/lib/multibeam_sonar/libmultibeam_sonar.so" test -f "$(ros2 pkg prefix multibeam_sonar_system)/lib/multibeam_sonar_system/libmultibeam_sonar_system.so" ``` -각 명령의 종료 상태가 0인지 확인하세요. 빌드·실행 로그에서 타깃 생략과 로드 오류를 확인하고, [멀티빔 소나](sonar.md)에서 영상·원시 데이터를 받습니다. +각 명령의 종료 상태가 0이어야 합니다. [소나 예제](sonar.md)를 실행하여 로그의 계산 방식을 확인하고 영상·원시 데이터를 수신하세요. 패키지 경로나 라이브러리 파일의 존재만으로 GPU 접근까지 확인할 수는 없습니다. + +## WGPU 어댑터 선택 + +Gazebo를 실행하는 터미널에서 다음 환경 변수를 설정합니다. + +| 변수 | 동작 | +| --- | --- | +| `WGPU_BACKEND` | GPU API 선택. Linux는 `vulkan`, macOS에서 직접 실행할 때는 `metal` 등 | +| `WGPU_ADAPTER_NAME` | 대소문자 구분 없이 값이 이름에 포함된 첫 번째 어댑터 선택 | +| `WGPU_POWER_PREF` | `high`·`low`·`none` 전력 선호 설정. 기본값은 `high` | + +예를 들어 이름에 `RTX`가 포함된 Vulkan 어댑터를 선택합니다. + +```bash +WGPU_BACKEND=vulkan WGPU_ADAPTER_NAME=RTX WGPU_POWER_PREF=high \ +ros2 launch posim_multibeam_sonar_demo multibeam_sonar_demo.launch.py \ + compute_backend:=wgpu +``` + +컴퓨터에서 사용 가능한 어댑터 이름을 지정하세요. 일치하는 이름이 없으면 로그에 사용 가능한 이름을 출력합니다. 선택된 WGPU 어댑터는 같은 Gazebo 프로세스의 소나 센서들이 공유합니다. 어댑터를 바꾸려면 새 프로세스를 시작하세요. + +## Apple Silicon + +macOS에서 직접 실행하려면 Gazebo Jetty·Ogre2·ROS–Gazebo 브리지가 포함된 ROS 2 작업 공간의 환경을 불러옵니다. 위 안내에 따라 Rust·Cargo와 소나 패키지를 빌드하세요. WGPU 계산은 Metal을 통해 Apple GPU를 사용합니다. + +```bash +WGPU_BACKEND=metal \ +ros2 launch posim_multibeam_sonar_demo multibeam_sonar_demo.launch.py \ + compute_backend:=wgpu +``` + +Metal은 음향 계산을 담당하며, Gazebo의 렌더링 환경도 필요합니다. CPU에서 음향을 계산하려면 `compute_backend:=cpu`를 사용하세요. CUDA에는 호환 NVIDIA 환경이 필요합니다. + +Apple Silicon 호스트의 Linux ARM64 컨테이너는 [Docker 설치와 가상 디스플레이 렌더링](docker.md)을 따르세요. 컨테이너 안에서 `POSIM_SONAR_COMPUTE_BACKEND=cpu`를 지정하여 [서버 모드 소나 명령](sonar.md)을 실행합니다. Linux 컨테이너에서 WGPU를 사용하려면 Vulkan 어댑터가 필요합니다. 컨테이너의 계산 API는 macOS에서 직접 실행할 때의 Metal과 별도로 구성합니다. + +## 시작 오류 해결 + +| 증상 | 조치 | +| --- | --- | +| `wgpu_vendorConfig.cmake`를 찾지 못함 | `wgpu_vendor`를 빌드하고 `install/setup.bash`를 불러온 뒤 나머지 패키지 빌드 | +| WGPU에서 사용 가능한 어댑터가 없다고 표시 | GPU 드라이버·Linux의 `vulkaninfo --summary`·어댑터 선택 변수 확인. GPU 계산을 사용할 수 없으면 `auto`나 `cpu` 선택 | +| CUDA를 직접 선택했을 때 실패 | CUDA 빌드 지원·`nvidia-smi`·런타임 라이브러리·컨테이너 GPU 접근 확인 | +| 알 수 없는 계산 방식이라고 표시 | `auto`·`wgpu`·`cuda`·`cpu` 중 선택 | +| 같은 이름의 월드가 이미 실행 중이라고 표시 | 기존 POSIM Launch를 Ctrl+C로 종료한 뒤 다시 실행 | +| 계산 방식은 초기화되지만 영상이 오지 않음 | 월드의 일시 정지를 해제하고 렌더러와 [소나 토픽](sonar.md) 확인 | ## 한 번에 하나씩 조절 | 설정 | 확인할 영향 | | --- | --- | -| 수평 빔·수직 ray 수 | 각도 샘플링과 계산량 | +| 수평 빔·수직 광선 수 | 각도 샘플링과 계산량 | | 거리·`maxDistance` | 관측 범위와 거리 처리 | | `raySkips` | 샘플링과 계산량의 균형 | | 센서 `update_rate` | 요청하는 센서 처리율 | | `writeLog`, `debugFlag` | 디스크·콘솔 출력 부하 | | 센서 gain | 영상 표현. 물리적 반사 강도와 구분 | -같은 장면·자세에서 프레임 처리 시간, 수신 토픽 주기, Gazebo 실시간 비율, CPU·GPU와 메모리 사용량을 기록합니다. 물리 계산, 렌더링, CUDA 계산, 발행과 기록 비용을 구분하세요. 측정 전에 렌더러를 충분히 실행합니다. +같은 장면·자세에서 프레임 처리 시간, 수신 토픽 주기, Gazebo 실시간 비율과 CPU·GPU 사용량을 기록합니다. 물리 계산·렌더링·음향 계산·발행·기록 비용을 구분하세요. 렌더러를 충분히 실행한 뒤 측정합니다. 설정된 30 Hz는 요청 주기이며, 실제 수신 처리율은 장면·센서 설정·하드웨어에 따라 달라집니다. -계산 커널을 변경할 때는 기준 데이터를 보관합니다. 화면뿐 아니라 강도·거리 결과를 비교하세요. 관련 연구와 출처는 [인용 안내](citation.md)에 있습니다. +계산 방식이나 센서 설정을 비교할 때는 원시 강도·거리 데이터를 사용하고 gain·표시 설정을 일정하게 유지하세요. 연구 출처와 저자 표시는 [인용 안내](citation.md)에 있습니다. ## GPU 계산 흐름
소나 계산의 GPU 처리 단계
소나 계산의 GPU 처리 단계
-소나 계산은 산란점 합산, 창 함수·보정과 FFT 단계로 구성됩니다. 프로파일링할 때 각 단계의 실행 시간과 GPU 메모리 전송을 구분하세요. 그림 출처: POSIM Notion Wiki. +소나 계산은 후방 산란·빔 보정·거리 압축으로 구성됩니다. 프로파일링할 때 계산 시간과 메모리 전송을 구분하세요. 그림 출처: POSIM Notion Wiki. diff --git a/website/content/ko/sonar.md b/website/content/ko/sonar.md index dddb508e..e6d3098c 100644 --- a/website/content/ko/sonar.md +++ b/website/content/ko/sonar.md @@ -1,22 +1,41 @@ # 멀티빔 소나 -POSIM 멀티빔 소나는 ray 기반 점 산란 모델로 강도·거리 데이터와 소나 영상을 생성합니다. 실행 전에 [NVIDIA/CUDA 환경](sonar-tuning.md)을 준비하세요. +POSIM 멀티빔 소나는 광선 기반 점 산란 모델로 강도·거리 데이터, 소나 영상과 점군을 생성합니다. 실행할 때 WGPU·CUDA·CPU 계산 방식을 선택합니다. 먼저 [소나 빌드와 렌더링 환경](sonar-tuning.md)을 준비하세요. ## 소나 예제 실행 +POSIM 작업 공간의 환경을 불러온 터미널에서 WGPU로 BlueView P900 예제를 실행합니다. + ```bash -ros2 launch posim_demos posim_sensor.launch.py \ - namespace:=blueview_p900 world_name:=posim_multibeam_sonar \ - paused:=false x:=4 z:=2.0 yaw:=3.14 gui:=false headless:=true +ros2 launch posim_multibeam_sonar_demo multibeam_sonar_demo.launch.py \ + compute_backend:=wgpu ``` -`posim_multibeam_sonar_demo` 패키지의 별도 예제도 있습니다. +예제는 Gazebo와 RViz를 엽니다. RViz에서 센서의 점군을 확인할 수 있습니다. RViz 창을 열지 않으려면 `rviz:=false`를 지정하세요. Launch는 하나씩 실행하고, 다음 Launch를 시작하기 전에 Ctrl+C로 종료합니다. + +## 계산 방식 선택 + +예제를 실행할 때 `compute_backend`를 지정합니다. 설치된 플러그인은 다시 컴파일하지 않고 다음 방식을 선택할 수 있습니다. + +| 값 | 선택 방식과 요구 사항 | +| --- | --- | +| `auto` (기본값) | CUDA, WGPU, CPU 순서로 사용 가능한 방식을 선택 | +| `wgpu` | WGPU 어댑터 사용. Linux는 Vulkan, macOS에서 직접 실행할 때는 Metal | +| `cuda` | NVIDIA GPU 사용. 플러그인의 CUDA 지원과 정상 작동하는 드라이버·런타임 필요 | +| `cpu` | CPU에서 음향 반사값 계산. Gazebo 렌더링도 필요 | + +실행 로그의 `Using sonar compute backend:`와 선택된 WGPU 어댑터를 확인하세요. `wgpu`나 `cuda`를 직접 지정하면 해당 방식이 초기화되어야 합니다. 사용할 수 없는 방식이나 잘못된 이름을 지정하면 해당 소나 센서를 비활성화하고 오류를 출력합니다. 자동 선택을 허용하려면 `auto`, CPU 계산을 직접 선택하려면 `cpu`를 사용하세요. + +일반 센서 Launch에서는 실행 환경의 `POSIM_SONAR_COMPUTE_BACKEND`를 설정합니다. ```bash -ros2 launch posim_multibeam_sonar_demo multibeam_sonar_demo.launch.py +POSIM_SONAR_COMPUTE_BACKEND=cpu \ +ros2 launch posim_demos posim_sensor.launch.py \ + namespace:=blueview_p900 world_name:=posim_multibeam_sonar \ + paused:=false x:=4 z:=2.0 yaw:=3.14 gui:=false headless:=true ``` -Launch는 하나씩 실행합니다. 로봇형 예제는 적절한 소나 월드에서 `bluerov2_heavy_multibeam_sonar`를 선택합니다. +서버 모드에도 렌더링 환경이 필요합니다. 가상 디스플레이나 GPU 연결은 [Docker 렌더링 안내](docker.md)를 따르세요. 로봇형 예제는 적절한 소나 월드에서 `bluerov2_heavy_multibeam_sonar`를 선택합니다. ## 출력 확인 @@ -24,9 +43,20 @@ Launch는 하나씩 실행합니다. 로봇형 예제는 적절한 소나 월드 ros2 topic list -t ros2 topic echo /sensor/multibeam_sonar/sonar_image \ sensor_msgs/msg/Image --once --no-arr +ros2 topic hz /sensor/multibeam_sonar/sonar_image ``` -RViz에서 소나 영상은 Image, `/sensor/multibeam_sonar/point_cloud`는 PointCloud2 디스플레이를 추가합니다. 원시 반사값은 `marine_acoustic_msgs`의 음향 메시지를 사용합니다. 정확한 자료형은 `ros2 topic list -t`로 찾습니다. +| 토픽 | ROS 2 메시지 | 용도 | +| --- | --- | --- | +| `/sensor/multibeam_sonar/sonar_image` | `sensor_msgs/msg/Image` | RViz의 Image 디스플레이에서 소나 강도 영상 표시 | +| `/sensor/multibeam_sonar/sonar_image_raw` | `marine_acoustic_msgs/msg/ProjectedSonarImage` | 거리·빔 정보를 포함한 음향 관측값 확인 | +| `/sensor/multibeam_sonar/point_cloud` | `sensor_msgs/msg/PointCloud2` | RViz의 PointCloud2 디스플레이에서 광선의 기하 정보 표시 | + +소나는 영상과 원시 음향 메시지를 ROS 2에 직접 발행합니다. 센서 설정의 브리지는 Gazebo 점군을 ROS 2로 전달합니다. 토픽 이름을 바꾼 경우 `ros2 topic list -t`로 확인하세요. + +## 센서 계산 흐름 + +Gazebo는 장면의 광선 거리와 표면 법선을 제공합니다. 소나 모델은 음향 후방 산란을 계산하고 빔 보정과 거리 압축을 수행합니다. 계산된 음향 데이터와 표시 영상은 같은 관측을 나타냅니다. 분석에는 원시 데이터를 사용하고, 영상과 함께 gain·표시 설정을 보관하세요. ## 센서 구성 @@ -35,7 +65,7 @@ RViz에서 소나 영상은 Image, `/sensor/multibeam_sonar/point_cloud`는 Poin | 설정 | 값 | | --- | --- | | 수평 빔 | 512개, 약 130° 시야 | -| 수직 ray | 300개, 약 12° 시야 | +| 수직 광선 | 300개, 약 12° 시야 | | 범위 | 0.1–10 m | | 주파수·대역폭 | 900 kHz / 29.9 kHz | | 음속 | 1500 m/s | @@ -45,13 +75,13 @@ RViz에서 소나 영상은 Image, `/sensor/multibeam_sonar/point_cloud`는 Poin ## 변경할 설정 -``에서 `sonarFreq`, `bandwidth`, `soundSpeed`, `sourceLevel`, `maxDistance`, `raySkips`, `sensorGain`과 출력 이름을 설정합니다. Ray 범위와 처리 거리를 맞추세요. `pointCloudTopicName`은 ROS–Gazebo 브리지 설정과 일치시키고 `frameName`에 소나 광학 좌표계를 지정합니다. +``에서 `sonarFreq`, `bandwidth`, `soundSpeed`, `sourceLevel`, `maxDistance`, `raySkips`, `sensorGain`과 출력 이름을 설정합니다. 광선 거리 범위와 처리 거리를 맞추세요. `pointCloudTopicName`은 ROS–Gazebo 브리지 설정과 일치시키고 `frameName`에 소나 광학 좌표계를 지정합니다. `writeLog`, `writeFrameInterval`, `debugFlag`로 진단 기록과 처리 시간을 확인할 수 있습니다. 기록 자체가 실험 목적이 아니라면 처리율 측정 시 기록을 끄세요. 반복 가능한 조절 순서는 [소나 빌드와 성능](sonar-tuning.md), 모델의 연구 출처는 [인용 안내](citation.md)에 있습니다. ## 소나 센서와 월드 추가하기 -`models/posim_sensor_models/description/blueview_p900/`을 새 센서 이름으로 복사합니다. 같은 이름의 `config/<센서 이름>/sensor_config.py`도 준비하고, SDF의 모델·링크·프레임과 브리지 토픽을 함께 맞춥니다. ``의 수평·수직 샘플 수와 각도, 거리 범위를 센서 사양에 맞게 설정합니다. +`models/posim_sensor_models/description/blueview_p900/`을 새 센서 이름으로 복사합니다. 같은 이름의 `config/<센서 이름>/sensor_config.py`도 준비하고, SDF의 모델·링크·프레임과 브리지 토픽을 함께 맞춥니다. ``의 수평 빔·수직 광선 수와 각도, 거리 범위를 센서 사양에 맞게 설정합니다. 월드에는 Gazebo 렌더링 센서 시스템과 소나 시스템을 포함합니다. 물리·장면·사용자 센서의 나머지 구성은 [소나 월드 SDF](https://github.com/IOES-Lab/POSIM/blob/main/models/posim_worlds/worlds/posim_multibeam_sonar.world)를 참고하세요. @@ -62,7 +92,7 @@ RViz에서 소나 영상은 Image, `/sensor/multibeam_sonar/point_cloud`는 Poin ``` -CUDA와 렌더링을 준비하고 작업 공간을 다시 빌드·source한 뒤, [센서 launch 인수](quickstart.md)를 사용해 새 descriptor를 선택합니다. ROS 점군 브리지는 `sensor_msgs/msg/PointCloud2`와 Gazebo `gz.msgs.PointCloudPacked`를 연결합니다. `pointCloudTopicName`과 `sensor_config.py`에 지정한 경로를 일치시키세요. +선택한 계산 방식과 렌더링을 준비하고 작업 공간을 다시 빌드하고 환경을 불러온 뒤, [센서 launch 인수](quickstart.md)를 사용해 새 센서 모델를 선택합니다. ROS 점군 브리지는 `sensor_msgs/msg/PointCloud2`와 Gazebo `gz.msgs.PointCloudPacked`를 연결합니다. `pointCloudTopicName`과 `sensor_config.py`에 지정한 경로를 일치시키세요. ## 원시 소나 데이터 기록하기 diff --git a/website/content/ko/troubleshooting.md b/website/content/ko/troubleshooting.md index eaacd03a..3275bf22 100644 --- a/website/content/ko/troubleshooting.md +++ b/website/content/ko/troubleshooting.md @@ -17,7 +17,7 @@ ros2 pkg prefix posim_demos ## 영상이 없거나 렌더링 초기화가 실패할 때 -영상 센서는 화면 없이 실행해도 렌더러가 필요합니다. 디스플레이·offscreen 설정을 확인하세요. [Docker 환경](docker.md)의 Xvfb 설정이나 [CUDA 라이브러리](sonar-tuning.md)를 확인합니다. +영상 센서는 화면 없이 실행해도 렌더러가 필요합니다. 디스플레이·offscreen 설정을 확인하세요. [Docker 환경](docker.md)의 Xvfb 설정이나 [소나 계산 환경](sonar-tuning.md)를 확인합니다. ## 메시지가 오지 않을 때 diff --git a/website/sources.json b/website/sources.json index d57c1cdd..0e5f7f81 100644 --- a/website/sources.json +++ b/website/sources.json @@ -1,5 +1,5 @@ { - "sourceRevision": "23bb608a9a89b9cda3472792316a81711f1ce73e", + "sourceRevision": "ad6165bd6c2b4e407f18ecb264e20bb76da8d4ca", "reviewedAt": "2026-10-09", "wiki": "https://caring-dibble-be5.notion.site/d24c9419989882cfa6498152ad4c840c?v=b42c941998988252a89588ad58631495&pvs=74", "notion": {