Photo-SLAM(3D高斯溅射SLAM)在 RDK S100 上的适配可行性探讨

各位好,

我正在研究将 Photo-SLAM(基于3D高斯溅射的实时SLAM系统)适配到 RDK S100 平台,遇到了一些架构层面的疑问,希望能和社区一起探讨。

核心矛盾:S100 没有 NVIDIA GPU,Photo-SLAM 严重依赖 CUDA

RDK S100 的硬件架构是 CPU (6×A78AE) + BPU Nash (80 TOPS) + MCU (4×R52+) + Mali-G78AE (仅 100 GFLOPS)。这意味着:

问题 详情
无 CUDA 设备 Mali-G78AE 是移动端 GPU,不支持 CUDA
BPU 不等于通用 NPU BPU Nash 专为 CNN/Transformer 推理优化,支持 160+ ONNX 算子,但不是通用 GPU 计算单元
光栅化管线完全无法运行 Photo-SLAM 的核心——可微高斯光栅化(tile-based α-blending + 基数排序 + 前缀和)是纯 CUDA kernel,无法在 Mali GPU 或 BPU 上直接执行
PyTorch CUDA 扩展失效 整个 src/rasterize_points.cu 的 ATen 绑定依赖 CUDA

所以这不是"简单移植几个算子"的问题,而是 整套渲染管线需要重写


一、项目简介

Photo-SLAM 是一个融合 ORB-SLAM3 前端与 3D Gaussian Splatting 后端的照片级真实感 SLAM 系统,GitHub: HuajianUP/Photo-SLAM。

其计算负载主要分为两部分:

模块 计算特征 当前实现
ORB-SLAM3 前端 特征提取、特征匹配、BA优化 纯 CPU,可移植
3DGS 渲染与训练 高斯光栅化、α-blending、协方差计算、SH→RGB、基数排序 纯 CUDA,不可移植

二、S100 平台适配的关键问题

我梳理了 3DGS 管线中 5 类 CUDA 计算任务在 S100 上的迁移可行性:

1. 高斯光栅化核心(前向渲染 + 反向传播)

  • 含 tile-based α-blending、cub::DeviceRadixSortcub::DeviceScan::InclusiveSum
  • 这是最大的计算瓶颈,也是最难迁移的部分
  • :red_question_mark: 问题:S100 的 Mali-G78AE GPU 能否通过 OpenCL/Vulkan Compute 实现等价的 tile-based 渲染?BPU 是否支持自定义算子注册来实现 α-blending 这类非标准算子?

2. 预处理算子(3D协方差计算 + SH→RGB 颜色转换)

  • 本质是逐高斯的矩阵运算,计算模式为 3×3 矩阵乘法和多项式求值
  • 理论上可以转为 ONNX 算子组合在 BPU 上推理
  • :red_question_mark: 问题:S100 的 BPU 工具链是否支持将自定义的逐元素矩阵运算(如3×3协方差计算)编译为 BPU 算子?160+ ONNX 算子中是否覆盖了 MatMul、ReduceSum、Where 等基础算子?

3. 点云操作与深度反投影

  • 纯 4×4 矩阵变换和逐元素乘加
  • 迁移难度低,CPU 即可完成
  • :white_check_mark: 这部分预计无障碍

4. KNN(Morton码排序 + 桶式最近邻搜索)

  • 依赖 CUB 的基数排序和 Thrust
  • :red_question_mark: 问题:S100 是否提供高性能排序库?BPU 或 CPU 上的排序性能如何?

5. PyTorch CUDA 扩展层

  • 当前通过 ATen C++ API 绑定 CUDA kernel
  • :red_question_mark: 问题:S100 是否支持 PyTorch 的运行?如果支持,是否有对应的后端(如 ARM PyTorch + OpenCL 后端)?还是需要切换到 ONNX Runtime / ncnn 等推理框架?

三、我目前的思路

初步考虑两种技术路线:

路线A:CPU + BPU 异构

  • ORB-SLAM3 前端在 CPU 上运行(A78AE 性能足够)
  • 3DGS 预处理算子(协方差、SH)转为 ONNX 模型在 BPU 上推理
  • 光栅化渲染在 CPU 上用多线程实现(性能会是瓶颈)
  • 放弃训练,只做推理渲染

路线B:Mali GPU OpenCL + BPU 混合

  • 预处理在 BPU 上
  • 光栅化通过 OpenCL/Vulkan Compute 在 Mali-G78AE 上实现
  • 需要将整个 CUDA 光栅化管线重写为 OpenCL

:red_question_mark: 想请教社区:

  1. Mali-G78AE 在 S100 上的 OpenCL/Vulkan Compute 支持情况如何?有无开发文档和示例?
  2. BPU 自定义算子的开发流程是什么?是否有教程?限制有哪些?
  3. S100 上是否有 PyTorch 的适配方案?或者推荐用什么推理框架替代?
  4. 是否有在 S100 上跑过类似的 SLAM + 神经渲染系统的先例?

四、环境信息

  • 目标平台:RDK S100 (80 TOPS, 12GB)
  • 源项目:Photo-SLAM (CUDA 11.8+, PyTorch 2.0+)
  • 我的开发背景:有 SLAM 和 3DGS 基础,但对 BPU 工具链不熟悉

感谢任何建议和指引!:handshake:

结论:Photo-SLAM 在 S100 上的适配是高难度但可行的,需要完全重构渲染管线

你的分析非常准确,这确实不是简单移植,而是架构级重写。基于社区现有资源和 S100 的硬件特性,我给你以下结论和建议:


一、核心问题解答

:one: Mali-G78AE 的 OpenCL/Vulkan Compute 支持情况

结论:技术上可用,但官方不支持

  • :white_check_mark: OpenCL 驱动存在(可通过 clinfo 验证),但:
    • :cross_mark: 无官方文档/教程
    • :cross_mark: 无技术支持保障
    • :warning: 性能和稳定性未验证
  • :pushpin: 官方建议:加速计算优先使用 BPU,GPU 仅用于图形渲染(OpenGL ES)

参考帖子请问 RDKx5 支持使用 opencl 调用 GPU 进行并行运算吗


:two: BPU 自定义算子开发流程

结论:支持自定义算子,但有约束

开发流程

ONNX 模型 → hb_mapper 检查算子支持 → 不支持则需自定义 → 编译为.hbm

关键约束

  • BPU Nash 支持 160+ ONNX 算子,但主要是 CNN/Transformer 类
  • 非标准算子(如 α-blending、tile-based 渲染)需要:
    1. 检查是否可用基础算子组合实现(MatMul、ReduceSum、Where 等)
    2. 如无法组合,需回退到 CPU 实现
    3. 或通过 quant_config 配置算子精度间接控制执行位置

参考资源


:three: S100 上的 PyTorch/推理框架方案

结论:不支持 PyTorch CUDA,需切换到 ONNX Runtime / HB DNN

框架 支持情况 建议
PyTorch CUDA :cross_mark: 不支持 放弃
PyTorch ARM (CPU) :white_check_mark: 可用 仅适合验证,性能差
ONNX Runtime :white_check_mark: 推荐 CPU 推理
HB DNN (BPU) :white_check_mark: 强烈推荐 BPU 加速推理
ncnn :warning: 可用 需自行移植

推荐路径

PyTorch 模型 → 导出 ONNX → hb_mapper 转换 → .hbm 格式 → HB DNN 推理

:four: S100 上 SLAM + 神经渲染的先例

结论:无 3DGS 先例,但有 SLAM 部署经验

  • :white_check_mark: ORB-SLAM3 已官方适配
  • :white_check_mark: 深度学习 SLAM 前端(如 Vista-SLAM)有用户成功在 CPU 上跑通,BPU 加速需手动部署
  • :cross_mark: 3DGS/神经渲染 暂无公开案例

参考帖子S100 上有什么适配的单目视觉 SLAM


二、技术路线建议

基于你的分析,我推荐 路线 A 的改进版

:bullseye: 推荐方案:CPU + BPU 异构(推理优先)

┌─────────────────────────────────────────────────────────┐
│ 模块                 │ 执行单元  │ 实现方式           │
├─────────────────────────────────────────────────────────┤
│ ORB-SLAM3 前端       │ CPU       │ 直接移植           │
│ 3DGS 预处理           │ BPU       │ ONNX → HBIR        │
│   - 协方差计算        │           │                    │
│   - SH→RGB 转换       │           │                    │
│ 高斯光栅化           │ CPU       │ 多线程重写         │
│ KNN 排序             │ CPU       │ 使用 STL/并行排序   │
│ 训练模块             │ ❌ 放弃    │ 仅做推理渲染        │
└─────────────────────────────────────────────────────────┘

为什么不选路线 B(Mali GPU OpenCL)?

  1. 无官方支持,踩坑成本高
  2. Mali-G78AE 仅 100 GFLOPS,性能可能不如 CPU 多线程
  3. 需要完全重写 CUDA 光栅化管线,工作量巨大

三、具体实施步骤

步骤 1:验证 CPU 版本可行性

# 1. 在 S100 上直接编译 Photo-SLAM(CPU 模式)
git clone https://github.com/HuajianUP/Photo-SLAM.git
cd Photo-SLAM
mkdir build && cd build

# 2. 修改 CMakeLists.txt,禁用 CUDA
# 注释掉所有 find_package(CUDA) 和 CUDA 相关编译选项

# 3. 用 CPU 实现光栅化(参考开源 CPU 实现)
# 可参考:https://github.com/graphdeco-inria/gaussian-splatting 的 CPU 分支

# 4. 测试帧率
./photo_slam --sequence <your_sequence>

预期结果:帧率可能只有 2-5 FPS,用于验证算法正确性


步骤 2:BPU 加速预处理算子

# 1. 导出预处理部分为 ONNX
# 将协方差计算、SH→RGB 转换为独立 ONNX 模型

# 2. 使用 hb_mapper 转换
hb_mapper toolkit \
  --model covariance_sh.onnx \
  --output output_dir \
  --march j6em  # S100 对应 J6 架构

# 3. 检查算子是否落在 BPU 上
# 编译日志会显示每个算子的执行位置

# 4. C++ 集成
# 使用 HB DNN Library 加载.hbm 文件
#include "hb_dnn.h"
// 参考:https://developer.d-robotics.cc/

参考文档RDK 模型部署指南


步骤 3:CPU 多线程光栅化优化

// 使用 OpenMP 并行化 tile-based 渲染
#pragma omp parallel for
for (int tile_y = 0; tile_y < num_tiles_y; tile_y++) {
    for (int tile_x = 0; tile_x < num_tiles_x; tile_x++) {
        render_tile(tile_x, tile_y, gaussians, image_buffer);
    }
}

// 使用 SIMD 优化(S100 支持 NEON)
// 编译时添加:-mfpu=neon -mfloat-abi=softfp

四、关键风险与应对

风险 影响 应对方案
光栅化 CPU 性能不足 帧率<5 FPS 降低高斯数量/分辨率,或放弃实时性
BPU 不支持关键算子 预处理无法加速 回退到 CPU,或简化算子
内存占用过高 12GB 不够用 限制高斯数量,使用内存池
无社区案例参考 踩坑成本高 建议将问题沉淀到社区,吸引官方关注

五、社区行动建议

  1. 将本帖问题补充到原帖,标记为技术探讨,吸引官方技术团队关注
  2. 如成功实现 CPU 版本,建议在社区分享经验(目前无 3DGS 案例)
  3. 关注官方动态:S100 作为新平台,后续可能提供更多算子支持

六、快速验证命令

# 1. 检查 OpenCL 可用性(不推荐依赖)
clinfo | grep "Platform #0"

# 2. 检查 BPU 工具链
which hb_mapper
hb_mapper --version

# 3. 查看 S100 算子支持列表
# https://toolchain.d-robotics.cc/guide/appendix/supported_op_list/operator_support/torch_operator_support_j6em.html

# 4. 测试 CPU 多线程性能
nproc  # 查看 CPU 核心数(应为 6)

总结:Photo-SLAM 在 S100 上的适配是可行的,但需要放弃训练、仅做推理,且光栅化部分需 CPU 重写。建议先用 CPU 跑通验证,再逐步用 BPU 加速预处理算子。如有具体实现问题,欢迎继续在社区讨论!:rocket:

建议放弃训练,训练放置到对应PC GPU进行,BPU可用于推理