wyj/GR-raytracing
1
0 Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Compare commits

base: wyj/GR-raytracing:265d7b95d5b916ec4768d174d73ad963a43f0088
wyj/GR-raytracing:main
..
compare: wyj/GR-raytracing:ab7d879f6cdb8a0f5d2d8ec4a4c2b40c0e45e070
wyj/GR-raytracing:main
2 Commits
265d7b95d5 .. ab7d879f6c
Author SHA1 Message Date
wyj ab7d879f6c Doc: record full-sky CPU and HIP benchmark comparison 
Preserve user-provided complete CPU and HIP logs and associate the optimized implementation with e1ec480669. Record matching event counts, user-confirmed bit-identical PNG output, and the 2.77x workflow speedup with tracing/import differences noted.

Mark the pre-artifact-fix CPU benchmark as historical and link the new paired result.
 2026-09-06 21:40:51 -04:00
wyj e1ec480669 HIP: accelerate PSF accumulation and restore parallel producers 
Cooperate across 32 lanes per PSF and use two completion-protected staging slots with complete batch timing. Restore coarse OpenMP event production while serializing shared GPU submissions and direct fallback boundaries.

Add bounded benchmarks, streaming and renderer regressions, and preserve validation evidence and ownership documentation.
 2026-09-06 21:38:53 -04:00
49 changed files with 10686 additions and 95 deletions
Expand all files Collapse all files

No files matched your search

Makefile
+9 -4
View File
@@ -41,9 +41,10 @@ SCHWARZSCHILD_TEST_TARGET := $(BUILD_DIR)/test_schwarzschild
OBSERVER_TRACK_TEST_TARGET := $(BUILD_DIR)/test_observer_track
CATALOG_PREFETCH_TEST_TARGET := $(BUILD_DIR)/test_catalog_prefetch
HIP_PSF_TEST_TARGET := $(BUILD_DIR)/test_hip_psf
HIP_PSF_BENCH_TARGET := $(BUILD_DIR)/benchmark_hip_psf
CAMERA_TEST_TARGETS := $(BUILD_DIR)/test_observer_minkowski $(BUILD_DIR)/test_observer_schwarzschild
.PHONY: all backend clean run test hip-psf-test minkowski schwarzschild FORCE
.PHONY: all backend clean run test hip-psf-test hip-psf-bench minkowski schwarzschild FORCE
ifneq ($(filter 0 1,$(PSF_EVENT_SINK)),$(PSF_EVENT_SINK))
$(error Unknown PSF_EVENT_SINK '$(PSF_EVENT_SINK)'; choose 0 or 1)
@@ -129,12 +130,16 @@ $(OBJECT_DIR)/$(HDR_BUILD_TAG)/src/hip_psf.o: src/hip_psf.hip src/hip_psf.h src/
# It requires PSF_BACKEND=hip and an actual HIP GPU agent at execution time.
ifeq ($(PSF_BACKEND),hip)
hip-psf-test: $(HIP_PSF_TEST_TARGET)
hip-psf-bench: $(HIP_PSF_BENCH_TARGET)
$(HIP_PSF_BENCH_TARGET): tests/benchmark_hip_psf.hip $(HIP_PSF_OBJECT) $(OBJECT_DIR)/$(HDR_BUILD_TAG)/src/optics.o | $(BUILD_DIR)
$(HIPCC) $(HIP_CXXFLAGS) $(OPENMP_FLAGS) -Isrc -x hip $< -x none $(filter-out $<,$^) $(LDLIBS) $(HDR_LDLIBS) -o $@
$(HIP_PSF_TEST_TARGET): tests/test_hip_psf.hip $(HIP_PSF_OBJECT) $(OBJECT_DIR)/$(HDR_BUILD_TAG)/src/optics.o | $(BUILD_DIR)
$(HIPCC) $(HIP_CXXFLAGS) $(OPENMP_FLAGS) -Isrc -x hip $< -x none $(filter-out $<,$^) $(LDLIBS) -o $@
$(HIPCC) $(HIP_CXXFLAGS) $(OPENMP_FLAGS) -Isrc -x hip $< -x none $(filter-out $<,$^) $(LDLIBS) $(HDR_LDLIBS) -o $@
else
hip-psf-test:
$(error hip-psf-test requires PSF_BACKEND=hip)
hip-psf-test hip-psf-bench:
$(error HIP PSF checks require PSF_BACKEND=hip)
endif
run: $(TARGET)
README.md
+3
View File
@@ -33,6 +33,9 @@ adaptive lens meshes, and reusable lens-map files. It is written primarily in
C with OpenMP CPU parallelism; an optional HIP backend accelerates PSF
accumulation.
HIP retains parallel CPU catalog mapping and uses bounded, completion-protected
event uploads. See [HIP configuration and bounded performance checks](build.md#optional-hip-psf-acceleration).
The [Nmesh](https://github.com/nmeshsource/nmesh) numerical-spacetime backend and BBH rendering are still planned.
The current scope is black-hole capture and distant stellar backgrounds;
local matter emission, accretion disks, and plasma are outside this stage.
README.zh-CN.md
+2
View File
@@ -19,6 +19,8 @@
当前实现支持解析 **Minkowski** 与 **Schwarzschild** 时空、单张图像、基于观测者轨迹的图像序列、自适应透镜网格，以及可复用的透镜映射文件。程序主要使用 C 编写，通过 OpenMP 实现 CPU 并行；可选的 HIP 后端用于加速 PSF 累积。
HIP 保留并行 CPU catalog 映射，并使用由完成事件保护的有界上传缓冲。配置与小规模性能检查见 [HIP 构建说明](build.md#optional-hip-psf-acceleration)。
[Nmesh](https://github.com/nmeshsource/nmesh) 数值时空后端和双黑洞（BBH）渲染仍在规划中。当前范围为黑洞捕获与远处恒星背景，暂不包含局域物质辐射、吸积盘或等离子体。架构与开发路线参见[设计文档](nr_spacetime_movie_renderer_design.md)。
## 编译
benchmarks/2mass_galactic_center_blackhole.md
+4
View File
@@ -1,3 +1,7 @@
> 历史记录：此运行早于 `265d7b95d5b916ec4768d174d73ad963a43f0088` 的伪像修复。
> 当前 CPU/HIP 配对基准及优化提交 hash 见 [完整 4K 新记录](2mass_galactic_center_cpu_hip_2026-09-07.md)。
> 下方原始命令与输出保留，不用于新版本计数对照。
Git hash: dd8cf29dac557c881e5df1db15a58fae15294967
```
benchmarks/2mass_galactic_center_cpu_hip_2026-09-07.md
+89
View File
@@ -0,0 +1,89 @@
# 2MASS 银河中心完整 4K：CPU / 优化 HIP（2026-09-07 归档）
## 代码版本与证据来源
HIP 优化提交：**`e1ec480669b0d24a0266ee135edf8a64a0d80ec9`**
（`HIP: accelerate PSF accumulation and restore parallel producers`）。
该提交包含 32-thread/event kernel、并行 CPU producer、有界双缓冲、全批次计时及回归测试。
HIP 完整运行发生在提交之前，对应此次提交的实现；运行日志本身没有嵌入 Git 或 binary hash。
用户指定本页 CPU 运行作为可比较基准，并说明
`265d7b95d5b916ec4768d174d73ad963a43f0088` 修复了此前的一些伪像。
因此，旧 `dd8cf29` 记录的 598,308,919 星像和 1,365 wing-clipped 不用于本次性能/正确性配对。
CPU 日志同样未嵌入 binary hash，不将代码版本说明冒充运行时二进制校验。
两份完整命令与原始终端输出均由用户提供：
- [CPU 原始日志](2mass_galactic_center_cpu_hip_2026-09-07/cpu.log)：生成并导出 lens map。
- [HIP 原始日志](2mass_galactic_center_cpu_hip_2026-09-07/hip.log)：导入同一 lens map；保留全部
worker 进度、每五秒 batch 进度、最终计数及 zsh `time` 输出，未删减成汇总。
用户确认：**两个 PNG 比特级完全一致。** 本次归档没有重跑渲染，也没有重新执行文件比较；
原始比较命令和两个 PNG 的 SHA256 未提供。不将 PNG 一致性扩大为完整任务的 linear HDR
逐样本一致性声明。此前有界样本的 HDR 验证另见
[有界排查记录](hip_psf_2026-09-07.md)。
## 工作负载与环境
- 主机沿用有界排查的 AMD Radeon RX 9070（gfx1201）与 HIP 7.2 环境；用户说明测试时段无其他
持续负载。完整日志没有独立输出驱动/编译器版本，前期环境证据见
[metadata.log](hip_psf_2026-09-07/metadata.log)。
- Release CPU/HIP 可执行文件，启用 HDR 输出；完整运行的实际编译命令未附在用户日志中。
- 3840×2160，45° FOV，RA 262.5° / Dec −30°，observer radius 100 M。
- CPU coarse cell 16，refine max level 4，Jacobian min 0.2；导出的 lens map 包含
65,817 vertices 和 130,882 triangles，HIP 直接导入它。
- 两次均请求并加载 64,440 tiles、343,363,141 stars，无 unavailable tile；catalog loaders=4。
- HIP 明确报告 16 producer workers。CPU 日志报告 1521% CPU，没有单独打印 worker 数。
- PSF FWHM 2.7 px、Moffat beta 4.5、relative tail 1e-8、min-Y=0、max-cache-psf-flux=1e8、
exposure=1e13；cache 为 64×64 phases、radius 47 px。
- 用户原始 CPU/HIP 命令使用相同 PNG 输出路径，且 CPU 命令的文件名也带 `full-gpu`；
原样保留此命名，不根据文件名判断 backend。未来复现应分别指定输出文件名，避免覆盖对照。
## 完整运行结果
| 项目 | CPU | 优化 HIP |
| --- | ---: | ---: |
| 总墙钟时间 | 1474.32 s（24:34.32） | 533.06 s（8:53.06） |
| user | 22234.08 s | 7733.61 s |
| system | 200.30 s | 61.18 s |
| CPU 百分比 | 1521% | 1462% |
| PSF cache build | 0.676 s | 0.691 s |
| catalog prefetch | 49.253 s | 48.201 s |
| 星像 / cached | 598264924 | 598264924 |
| wing-clipped | 0 | 0 |
| direct fallback | 0 | 0 |
| discarded below min-Y | 0 | 0 |
总墙钟时间比 `1474.32 / 533.06 = 2.76577`，耗时减少约 **63.84%**，节省 **941.26 秒**
（15 分 41.26 秒）。这是两次工作流的端到端时间比：CPU 包含 tracing 与 lens-map 写出，
HIP 导入文件；没有独立的 CPU PSF 阶段计时，不能把该比值称作严格隔离的 PSF backend 加速比。
HIP 全程诊断：
| 项目 | 结果 |
| --- | ---: |
| producer/splat wall | 481.155 s（frame 日志四舍五入为481.2 s） |
| GPU kernel | 476.906037 s |
| event upload | 1.504570 s |
| HDR download | 0.021966 s |
| batch / timed batch | 36526 / 36526 |
| summed worker generation | 3302.459 s |
| summed worker submission/fallback | 4395.356 s |
GPU kernel 时间约为 producer/splat wall 的 **99.12%**，kernel 吞吐约 **125.45 万事件/秒**。
上传占比很小，CPU producer 在中段能够持续供给 GPU；下一步性能工作应优先研究 kernel，
而不是仅增加 CPU worker 或上传带宽。这不单独证明 atomic contention 是 kernel 的主导成本。
多数进度行的 submitted−completed=32,768，等于两个 16,384-event staging slot，符合
有界队列与计时回收的设计。最后全部 36,526 批计时，没有沿用旧版前 64 批采样限制。
两个 worker 时间是线程墙钟之和，其中 submission/fallback 包含锁与设备等待：
`3302.459 + 4395.356 = 7697.815 ≈ 16 × 481.155`。不能将它们当作额外串行时间加到
GPU 或总墙钟时间上，1462% CPU 也不能全部解释为有效计算。
## 比较结论
本次完整任务确认了优化的实际收益，星像及 PSF 分类计数全部匹配，PNG 比特级一致由用户确认。
旧 CPU 文档中的计数差异属于不同修复版本，已不构成本次 CPU/HIP 配对的疑点。
原有完整日志、硬件/构建证据的限制与 tracing/import 差异均保留，不用小规模 kernel 加速比
替代这里的完整运行结果。
benchmarks/2mass_galactic_center_cpu_hip_2026-09-07/cpu.log
+23
View File
@@ -0,0 +1,23 @@
❯ time ./build/Release/schwarzschild_sky \
--psf-relative-tail 1e-8 \
--psf-min-y 0 \
--max-cache-psf-flux 1e8 \
--all-sky-catalog assets/2mass/processed/all_sky \
--refine-max-level 4 \
--refine-jacobian-min 0.2 \
--width 3840 --height 2160 \
--look-ra-deg 262.5 --look-dec-deg -30 \
--fov-deg 45 \
--coarse-cell-pixels 16 \
--exposure 1e13 \
--psf-fwhm-pixels 2.7 --psf-moffat-beta 4.5 \
--catalog-load-workers 4 \
--observer-radius 100 \
--hdr-output --lens-map-output output/lens/schwarzschild_galactic_center_R100_45deg_16_4_4k_.grlens \
--output output/imgs/2mass_galactic_center_schwarzschild_R100_4k_1e13_beta4.5_16_4-full-gpu.png
PSF cache ready: 64x64 phases, radius 47 px, relative tail 1e-08, tail abs 1e-06, boundary 1e-07, build 0.676 s
Wrote lens map: output/lens/schwarzschild_galactic_center_R100_45deg_16_4_4k_.grlens (65817 vertices, 130882 triangles)
Catalog prefetch: 64440 requested, 64440 newly loaded (343363141 stars), 0 unavailable in 49.253 s; 4 loader workers
Rendered 598264924 images from 343363141 catalog stars to output/imgs/2mass_galactic_center_schwarzschild_R100_4k_1e13_beta4.5_16_4-full-gpu.png (ok)
PSF splats: cached 598264924, cached wing-clipped 0, direct fallbacks 0, discarded below min-Y 0
./build/Release/schwarzschild_sky --psf-relative-tail 1e-8 --psf-min-y 0 1e8 22234.08s user 200.30s system 1521% cpu 24:34.32 total
benchmarks/2mass_galactic_center_cpu_hip_2026-09-07/hip.log
+315
View File
@@ -0,0 +1,315 @@
❯ time ./build/Release/schwarzschild_sky_hip --verbose \
--psf-relative-tail 1e-8 \
--psf-min-y 0 \
--max-cache-psf-flux 1e8 \
--all-sky-catalog assets/2mass/processed/all_sky \
--exposure 1e13 \
--psf-fwhm-pixels 2.7 --psf-moffat-beta 4.5 \
--catalog-load-workers 4 \
--hdr-output --lens-map-input output/lens/schwarzschild_galactic_center_R100_45deg_16_4_4k_.grlens \
--output output/imgs/2mass_galactic_center_schwarzschild_R100_4k_1e13_beta4.5_16_4-full-gpu.png
PSF cache ready: 64x64 phases, radius 47 px, relative tail 1e-08, tail abs 1e-06, boundary 1e-07, build 0.691 s
Frame 0: finding and prefetching catalog tiles...
Catalog prefetch: 64440 requested, 64440 newly loaded (343363141 stars), 0 unavailable in 48.201 s; 4 loader workers
Frame 0: catalog prefetch finished (64440 candidate tiles).
Frame 0: splatting 130882 lens triangles...
Frame 0: splat worker 13/16 started.
Frame 0: splat worker 1/16 started.
Frame 0: splat worker 10/16 started.
Frame 0: splat worker 4/16 started.
Frame 0: splat worker 8/16 started.
Frame 0: splat worker 11/16 started.
Frame 0: splat worker 3/16 started.
Frame 0: splat worker 15/16 started.
Frame 0: splat worker 14/16 started.
Frame 0: splat worker 6/16 started.
Frame 0: splat worker 5/16 started.
Frame 0: splat worker 9/16 started.
Frame 0: splat worker 16/16 started.
Frame 0: splat worker 7/16 started.
Frame 0: splat worker 12/16 started.
Frame 0: splat worker 2/16 started.
Frame 0: splat worker 2/16 reached 8 local triangles.
Frame 0: splat worker 12/16 reached 8 local triangles.
Frame 0: splat worker 10/16 reached 8 local triangles.
Frame 0: splat worker 1/16 reached 8 local triangles.
Frame 0: splat worker 11/16 reached 8 local triangles.
Frame 0: splat worker 5/16 reached 8 local triangles.
Frame 0: splat worker 4/16 reached 8 local triangles.
Frame 0: splat worker 6/16 reached 8 local triangles.
Frame 0: splat worker 9/16 reached 8 local triangles.
Frame 0: splat worker 3/16 reached 8 local triangles.
Frame 0: splat worker 16/16 reached 8 local triangles.
Frame 0: splat worker 7/16 reached 8 local triangles.
Frame 0: splat worker 15/16 reached 8 local triangles.
Frame 0: splat worker 8/16 reached 8 local triangles.
Frame 0: splat worker 13/16 reached 8 local triangles.
Frame 0: splat worker 14/16 reached 8 local triangles.
Frame 0: splat worker 6/16 reached 16 local triangles.
Frame 0: splat worker 10/16 reached 16 local triangles.
Frame 0: splat worker 1/16 reached 16 local triangles.
Frame 0: splat worker 4/16 reached 16 local triangles.
Frame 0: splat worker 16/16 reached 16 local triangles.
Frame 0: splat worker 12/16 reached 16 local triangles.
Frame 0: splat worker 13/16 reached 16 local triangles.
Frame 0: splat worker 5/16 reached 16 local triangles.
Frame 0: splat worker 14/16 reached 16 local triangles.
Frame 0: splat worker 11/16 reached 16 local triangles.
Frame 0: splat worker 7/16 reached 16 local triangles.
Frame 0: splat worker 15/16 reached 16 local triangles.
Frame 0: splat worker 8/16 reached 16 local triangles.
Frame 0: splat worker 2/16 reached 16 local triangles.
Frame 0: splat worker 3/16 reached 16 local triangles.
Frame 0: splat worker 9/16 reached 16 local triangles.
Frame 0: splat worker 12/16 reached 32 local triangles.
Frame 0: splat worker 13/16 reached 32 local triangles.
Frame 0: splat worker 1/16 reached 32 local triangles.
Frame 0: splat worker 4/16 reached 32 local triangles.
Frame 0: splat worker 6/16 reached 32 local triangles.
Frame 0: splat worker 12/16 reached 64 local triangles.
Frame 0: splat worker 13/16 reached 64 local triangles.
Frame 0: splat worker 14/16 reached 32 local triangles.
Frame 0: splat worker 5/16 reached 32 local triangles.
Frame 0: splat worker 10/16 reached 32 local triangles.
Frame 0: splat worker 15/16 reached 32 local triangles.
Frame 0: splat worker 2/16 reached 32 local triangles.
Frame 0: splat worker 7/16 reached 32 local triangles.
Frame 0: splat worker 11/16 reached 32 local triangles.
Frame 0: splat worker 16/16 reached 32 local triangles.
Frame 0: splat worker 3/16 reached 32 local triangles.
Frame 0: splat worker 9/16 reached 32 local triangles.
Frame 0: splat worker 8/16 reached 32 local triangles.
Frame 0: splat worker 1/16 reached 64 local triangles.
Frame 0: splat worker 4/16 reached 64 local triangles.
Frame 0: splat worker 6/16 reached 64 local triangles.
Frame 0: splat worker 14/16 reached 64 local triangles.
Frame 0: splat worker 9/16 reached 64 local triangles.
Frame 0: splat worker 3/16 reached 64 local triangles.
Frame 0: splat worker 16/16 reached 64 local triangles.
Frame 0: splat worker 7/16 reached 64 local triangles.
Frame 0: splat worker 8/16 reached 64 local triangles.
Frame 0: splat worker 11/16 reached 64 local triangles.
Frame 0: splat worker 12/16 reached 128 local triangles.
Frame 0: splat worker 5/16 reached 64 local triangles.
Frame 0: splat worker 13/16 reached 128 local triangles.
Frame 0: splat worker 1/16 reached 128 local triangles.
Frame 0: splat worker 15/16 reached 64 local triangles.
Frame 0: splat worker 10/16 reached 64 local triangles.
Frame 0: splat worker 2/16 reached 64 local triangles.
Frame 0: splat worker 4/16 reached 128 local triangles.
Frame 0: splat worker 9/16 reached 128 local triangles.
Frame 0: splat worker 3/16 reached 128 local triangles.
Frame 0: splat worker 8/16 reached 128 local triangles.
Frame 0: splat worker 6/16 reached 128 local triangles.
Frame 0: splat worker 11/16 reached 128 local triangles.
Frame 0: splat worker 14/16 reached 128 local triangles.
Frame 0: splat worker 16/16 reached 128 local triangles.
Frame 0: splat worker 5/16 reached 128 local triangles.
Frame 0: splat worker 10/16 reached 128 local triangles.
Frame 0: splat worker 7/16 reached 128 local triangles.
Frame 0: splat worker 2/16 reached 128 local triangles.
Frame 0: splat worker 15/16 reached 128 local triangles.
Frame 0: splat worker 13/16 reached 256 local triangles.
Frame 0: splat worker 12/16 reached 256 local triangles.
Frame 0: splat worker 1/16 reached 256 local triangles.
Frame 0: splat worker 4/16 reached 256 local triangles.
Frame 0: splat worker 3/16 reached 256 local triangles.
Frame 0: splat worker 14/16 reached 256 local triangles.
Frame 0: splat worker 8/16 reached 256 local triangles.
Frame 0: splat worker 9/16 reached 256 local triangles.
Frame 0: splat worker 10/16 reached 256 local triangles.
Frame 0: splat worker 5/16 reached 256 local triangles.
Frame 0: splat worker 7/16 reached 256 local triangles.
Frame 0: splat worker 6/16 reached 256 local triangles.
Frame 0: splat worker 16/16 reached 256 local triangles.
Frame 0: splat worker 15/16 reached 256 local triangles.
Frame 0: splat worker 11/16 reached 256 local triangles.
Frame 0: splat worker 2/16 reached 256 local triangles.
Frame 0: splat worker 4/16 reached 512 local triangles.
Frame 0: splat worker 12/16 reached 512 local triangles.
Frame 0: splat worker 13/16 reached 512 local triangles.
Frame 0: splat worker 9/16 reached 512 local triangles.
Frame 0: splat worker 10/16 reached 512 local triangles.
Frame 0: splat worker 8/16 reached 512 local triangles.
Frame 0: splat worker 3/16 reached 512 local triangles.
Frame 0: splat worker 1/16 reached 512 local triangles.
Frame 0: splat worker 14/16 reached 512 local triangles.
Frame 0: splat worker 6/16 reached 512 local triangles.
Frame 0: splat worker 15/16 reached 512 local triangles.
Frame 0: splat worker 5/16 reached 512 local triangles.
Frame 0: splat worker 11/16 reached 512 local triangles.
Frame 0: splat worker 16/16 reached 512 local triangles.
Frame 0: splat worker 7/16 reached 512 local triangles.
Frame 0: splat worker 2/16 reached 512 local triangles.
HIP PSF progress: 4816896 submitted / 4784128 completed events; 292 / 294 batches timed; kernel 4.083 s
Frame 0: splat worker 1/16 reached 1024 local triangles.
Frame 0: splat worker 4/16 reached 1024 local triangles.
Frame 0: splat worker 13/16 reached 1024 local triangles.
Frame 0: splat worker 12/16 reached 1024 local triangles.
Frame 0: splat worker 14/16 reached 1024 local triangles.
Frame 0: splat worker 8/16 reached 1024 local triangles.
Frame 0: splat worker 9/16 reached 1024 local triangles.
Frame 0: splat worker 6/16 reached 1024 local triangles.
Frame 0: splat worker 3/16 reached 1024 local triangles.
Frame 0: splat worker 10/16 reached 1024 local triangles.
Frame 0: splat worker 7/16 reached 1024 local triangles.
Frame 0: splat worker 15/16 reached 1024 local triangles.
Frame 0: splat worker 16/16 reached 1024 local triangles.
Frame 0: splat worker 2/16 reached 1024 local triangles.
HIP PSF progress: 11206656 submitted / 11173888 completed events; 682 / 684 batches timed; kernel 9.065 s
Frame 0: splat worker 11/16 reached 1024 local triangles.
Frame 0: splat worker 5/16 reached 1024 local triangles.
HIP PSF progress: 17547264 submitted / 17514496 completed events; 1069 / 1071 batches timed; kernel 14.057 s
HIP PSF progress: 23838720 submitted / 23805952 completed events; 1453 / 1455 batches timed; kernel 19.045 s
HIP PSF progress: 30081024 submitted / 30048256 completed events; 1834 / 1836 batches timed; kernel 24.036 s
HIP PSF progress: 36306944 submitted / 36274176 completed events; 2214 / 2216 batches timed; kernel 29.028 s
HIP PSF progress: 42483712 submitted / 42450944 completed events; 2591 / 2593 batches timed; kernel 34.020 s
HIP PSF progress: 48660480 submitted / 48627712 completed events; 2968 / 2970 batches timed; kernel 39.014 s
HIP PSF progress: 54755328 submitted / 54722560 completed events; 3340 / 3342 batches timed; kernel 43.998 s
HIP PSF progress: 60850176 submitted / 60817408 completed events; 3712 / 3714 batches timed; kernel 48.971 s
Frame 0: splat worker 14/16 reached 2048 local triangles.
Frame 0: splat worker 12/16 reached 2048 local triangles.
Frame 0: splat worker 4/16 reached 2048 local triangles.
Frame 0: splat worker 9/16 reached 2048 local triangles.
HIP PSF progress: 66977792 submitted / 66945024 completed events; 4086 / 4088 batches timed; kernel 53.963 s
Frame 0: splat worker 1/16 reached 2048 local triangles.
Frame 0: splat worker 3/16 reached 2048 local triangles.
Frame 0: splat worker 8/16 reached 2048 local triangles.
Frame 0: splat worker 16/16 reached 2048 local triangles.
Frame 0: splat worker 7/16 reached 2048 local triangles.
Frame 0: splat worker 13/16 reached 2048 local triangles.
Frame 0: splat worker 5/16 reached 2048 local triangles.
Frame 0: splat worker 15/16 reached 2048 local triangles.
Frame 0: splat worker 2/16 reached 2048 local triangles.
HIP PSF progress: 73138176 submitted / 73105408 completed events; 4462 / 4464 batches timed; kernel 58.945 s
Frame 0: splat worker 6/16 reached 2048 local triangles.
Frame 0: splat worker 10/16 reached 2048 local triangles.
Frame 0: splat worker 11/16 reached 2048 local triangles.
HIP PSF progress: 79216640 submitted / 79183872 completed events; 4833 / 4835 batches timed; kernel 63.931 s
HIP PSF progress: 85393408 submitted / 85360640 completed events; 5210 / 5212 batches timed; kernel 68.919 s
HIP PSF progress: 91537408 submitted / 91504640 completed events; 5585 / 5587 batches timed; kernel 73.903 s
HIP PSF progress: 97779712 submitted / 97746944 completed events; 5966 / 5968 batches timed; kernel 78.894 s
HIP PSF progress: 103890944 submitted / 103858176 completed events; 6339 / 6341 batches timed; kernel 83.877 s
HIP PSF progress: 110215168 submitted / 110182400 completed events; 6725 / 6727 batches timed; kernel 88.859 s
HIP PSF progress: 116506624 submitted / 116473856 completed events; 7109 / 7111 batches timed; kernel 93.844 s
HIP PSF progress: 123224064 submitted / 123191296 completed events; 7519 / 7521 batches timed; kernel 98.830 s
HIP PSF progress: 130007040 submitted / 129974272 completed events; 7933 / 7935 batches timed; kernel 103.809 s
HIP PSF progress: 136577024 submitted / 136544256 completed events; 8334 / 8336 batches timed; kernel 108.791 s
HIP PSF progress: 142950400 submitted / 142917632 completed events; 8723 / 8725 batches timed; kernel 113.777 s
HIP PSF progress: 149487616 submitted / 149454848 completed events; 9122 / 9124 batches timed; kernel 118.758 s
HIP PSF progress: 156352512 submitted / 156319744 completed events; 9541 / 9543 batches timed; kernel 123.748 s
HIP PSF progress: 163250176 submitted / 163217408 completed events; 9962 / 9964 batches timed; kernel 128.732 s
HIP PSF progress: 169689088 submitted / 169656320 completed events; 10355 / 10357 batches timed; kernel 133.705 s
HIP PSF progress: 176259072 submitted / 176226304 completed events; 10756 / 10758 batches timed; kernel 138.691 s
HIP PSF progress: 183222272 submitted / 183189504 completed events; 11181 / 11183 batches timed; kernel 143.674 s
HIP PSF progress: 190251008 submitted / 190218240 completed events; 11610 / 11612 batches timed; kernel 148.656 s
HIP PSF progress: 197181440 submitted / 197148672 completed events; 12033 / 12035 batches timed; kernel 153.645 s
HIP PSF progress: 203554816 submitted / 203522048 completed events; 12422 / 12424 batches timed; kernel 158.639 s
HIP PSF progress: 210452480 submitted / 210419712 completed events; 12843 / 12845 batches timed; kernel 163.627 s
HIP PSF progress: 217448448 submitted / 217415680 completed events; 13270 / 13272 batches timed; kernel 168.608 s
HIP PSF progress: 224083968 submitted / 224051200 completed events; 13675 / 13677 batches timed; kernel 173.601 s
HIP PSF progress: 230670336 submitted / 230637568 completed events; 14077 / 14079 batches timed; kernel 178.581 s
HIP PSF progress: 237600768 submitted / 237568000 completed events; 14500 / 14502 batches timed; kernel 183.570 s
HIP PSF progress: 244121600 submitted / 244088832 completed events; 14898 / 14900 batches timed; kernel 188.555 s
HIP PSF progress: 250937344 submitted / 250904576 completed events; 15314 / 15316 batches timed; kernel 193.537 s
HIP PSF progress: 257392640 submitted / 257359872 completed events; 15708 / 15710 batches timed; kernel 198.520 s
HIP PSF progress: 264028160 submitted / 263995392 completed events; 16113 / 16115 batches timed; kernel 203.508 s
HIP PSF progress: 270680064 submitted / 270647296 completed events; 16519 / 16521 batches timed; kernel 208.496 s
HIP PSF progress: 277184512 submitted / 277151744 completed events; 16916 / 16918 batches timed; kernel 213.465 s
HIP PSF progress: 283721728 submitted / 283688960 completed events; 17315 / 17317 batches timed; kernel 218.447 s
Frame 0: splat worker 14/16 reached 4096 local triangles.
HIP PSF progress: 290193408 submitted / 290160640 completed events; 17710 / 17712 batches timed; kernel 223.428 s
Frame 0: splat worker 16/16 reached 4096 local triangles.
HIP PSF progress: 296779776 submitted / 296747008 completed events; 18112 / 18114 batches timed; kernel 228.403 s
Frame 0: splat worker 2/16 reached 4096 local triangles.
HIP PSF progress: 303333376 submitted / 303300608 completed events; 18512 / 18514 batches timed; kernel 233.386 s
Frame 0: splat worker 13/16 reached 4096 local triangles.
Frame 0: splat worker 10/16 reached 4096 local triangles.
Frame 0: splat worker 7/16 reached 4096 local triangles.
HIP PSF progress: 309755904 submitted / 309723136 completed events; 18904 / 18906 batches timed; kernel 238.369 s
Frame 0: splat worker 4/16 reached 4096 local triangles.
Frame 0: splat worker 5/16 reached 4096 local triangles.
Frame 0: splat worker 8/16 reached 4096 local triangles.
HIP PSF progress: 315932672 submitted / 315899904 completed events; 19281 / 19283 batches timed; kernel 243.183 s
Frame 0: splat worker 3/16 reached 4096 local triangles.
Frame 0: splat worker 11/16 reached 4096 local triangles.
HIP PSF progress: 322338816 submitted / 322306048 completed events; 19672 / 19674 batches timed; kernel 248.175 s
Frame 0: splat worker 1/16 reached 4096 local triangles.
Frame 0: splat worker 6/16 reached 4096 local triangles.
Frame 0: splat worker 9/16 reached 4096 local triangles.
Frame 0: splat worker 15/16 reached 4096 local triangles.
Frame 0: splat worker 12/16 reached 4096 local triangles.
HIP PSF progress: 328597504 submitted / 328564736 completed events; 20054 / 20056 batches timed; kernel 253.161 s
HIP PSF progress: 334741504 submitted / 334708736 completed events; 20429 / 20431 batches timed; kernel 258.149 s
HIP PSF progress: 340819968 submitted / 340787200 completed events; 20800 / 20802 batches timed; kernel 263.138 s
HIP PSF progress: 346865664 submitted / 346832896 completed events; 21169 / 21171 batches timed; kernel 268.124 s
HIP PSF progress: 352894976 submitted / 352862208 completed events; 21537 / 21539 batches timed; kernel 273.097 s
HIP PSF progress: 358973440 submitted / 358940672 completed events; 21908 / 21910 batches timed; kernel 278.083 s
HIP PSF progress: 364986368 submitted / 364953600 completed events; 22275 / 22277 batches timed; kernel 283.074 s
HIP PSF progress: 370900992 submitted / 370868224 completed events; 22636 / 22638 batches timed; kernel 288.068 s
HIP PSF progress: 376946688 submitted / 376913920 completed events; 23005 / 23007 batches timed; kernel 293.054 s
HIP PSF progress: 382844928 submitted / 382812160 completed events; 23365 / 23367 batches timed; kernel 298.052 s
HIP PSF progress: 388923392 submitted / 388890624 completed events; 23736 / 23738 batches timed; kernel 303.033 s
HIP PSF progress: 394772480 submitted / 394739712 completed events; 24093 / 24095 batches timed; kernel 308.019 s
HIP PSF progress: 400621568 submitted / 400588800 completed events; 24450 / 24452 batches timed; kernel 313.006 s
HIP PSF progress: 406536192 submitted / 406503424 completed events; 24811 / 24813 batches timed; kernel 317.991 s
HIP PSF progress: 412499968 submitted / 412467200 completed events; 25175 / 25177 batches timed; kernel 322.978 s
HIP PSF progress: 418316288 submitted / 418283520 completed events; 25530 / 25532 batches timed; kernel 327.973 s
HIP PSF progress: 424148992 submitted / 424116224 completed events; 25886 / 25888 batches timed; kernel 332.964 s
HIP PSF progress: 429948928 submitted / 429916160 completed events; 26240 / 26242 batches timed; kernel 337.956 s
HIP PSF progress: 435961856 submitted / 435929088 completed events; 26607 / 26609 batches timed; kernel 342.950 s
HIP PSF progress: 441942016 submitted / 441909248 completed events; 26972 / 26974 batches timed; kernel 347.936 s
HIP PSF progress: 447741952 submitted / 447709184 completed events; 27326 / 27328 batches timed; kernel 352.922 s
HIP PSF progress: 453541888 submitted / 453509120 completed events; 27680 / 27682 batches timed; kernel 357.894 s
HIP PSF progress: 459390976 submitted / 459358208 completed events; 28037 / 28039 batches timed; kernel 362.875 s
HIP PSF progress: 465387520 submitted / 465354752 completed events; 28403 / 28405 batches timed; kernel 367.848 s
HIP PSF progress: 471400448 submitted / 471367680 completed events; 28770 / 28772 batches timed; kernel 372.843 s
HIP PSF progress: 476987392 submitted / 476954624 completed events; 29111 / 29113 batches timed; kernel 377.672 s
HIP PSF progress: 482525184 submitted / 482492416 completed events; 29449 / 29451 batches timed; kernel 382.488 s
HIP PSF progress: 488357888 submitted / 488325120 completed events; 29805 / 29807 batches timed; kernel 387.481 s
HIP PSF progress: 494190592 submitted / 494157824 completed events; 30161 / 30163 batches timed; kernel 392.467 s
HIP PSF progress: 500432896 submitted / 500400128 completed events; 30542 / 30544 batches timed; kernel 397.454 s
HIP PSF progress: 506445824 submitted / 506413056 completed events; 30909 / 30911 batches timed; kernel 402.437 s
HIP PSF progress: 512344064 submitted / 512311296 completed events; 31269 / 31271 batches timed; kernel 407.423 s
HIP PSF progress: 518455296 submitted / 518422528 completed events; 31642 / 31644 batches timed; kernel 412.408 s
HIP PSF progress: 524484608 submitted / 524451840 completed events; 32010 / 32012 batches timed; kernel 417.391 s
HIP PSF progress: 530628608 submitted / 530595840 completed events; 32385 / 32387 batches timed; kernel 422.377 s
HIP PSF progress: 536739840 submitted / 536707072 completed events; 32758 / 32760 batches timed; kernel 427.370 s
HIP PSF progress: 542867456 submitted / 542834688 completed events; 33132 / 33134 batches timed; kernel 432.358 s
HIP PSF progress: 549044224 submitted / 549011456 completed events; 33509 / 33511 batches timed; kernel 437.340 s
HIP PSF progress: 555220992 submitted / 555188224 completed events; 33886 / 33888 batches timed; kernel 442.323 s
HIP PSF progress: 561397760 submitted / 561364992 completed events; 34263 / 34265 batches timed; kernel 447.309 s
HIP PSF progress: 567607296 submitted / 567574528 completed events; 34642 / 34644 batches timed; kernel 452.300 s
HIP PSF progress: 573849600 submitted / 573816832 completed events; 35023 / 35025 batches timed; kernel 457.284 s
HIP PSF progress: 580124672 submitted / 580091904 completed events; 35406 / 35408 batches timed; kernel 462.268 s
HIP PSF progress: 586432512 submitted / 586399744 completed events; 35791 / 35793 batches timed; kernel 467.253 s
HIP PSF progress: 592773120 submitted / 592740352 completed events; 36178 / 36180 batches timed; kernel 472.247 s
Frame 0: splat worker 14/16 reached 8192 local triangles.
Frame 0: splat worker 16/16 reached 8192 local triangles.
Frame 0: splat worker 2/16 reached 8192 local triangles.
Frame 0: splat worker 13/16 reached 8192 local triangles.
Frame 0: splat worker 10/16 reached 8192 local triangles.
HIP PSF progress: 597475328 submitted / 597442560 completed events; 36465 / 36467 batches timed; kernel 476.352 s
Frame 0: splat worker 11/16 reached 8192 local triangles.
Frame 0: splat worker 3/16 finished after 8097 local triangles.
Frame 0: splat worker 12/16 finished after 7940 local triangles.
Frame 0: splat worker 13/16 finished after 8379 local triangles.
Frame 0: splat worker 1/16 finished after 8046 local triangles.
Frame 0: splat worker 16/16 finished after 8763 local triangles.
Frame 0: splat worker 9/16 finished after 7982 local triangles.
Frame 0: splat worker 11/16 finished after 8200 local triangles.
Frame 0: splat worker 8/16 finished after 8121 local triangles.
Frame 0: splat worker 14/16 finished after 8831 local triangles.
Frame 0: splat worker 4/16 finished after 7936 local triangles.
Frame 0: splat worker 6/16 finished after 8100 local triangles.
Frame 0: splat worker 5/16 finished after 7905 local triangles.
Frame 0: splat worker 15/16 finished after 7625 local triangles.
Frame 0: splat worker 7/16 finished after 8084 local triangles.
Frame 0: splat worker 10/16 finished after 8367 local triangles.
Frame 0: splat worker 2/16 finished after 8506 local triangles.
HIP producers: 16 workers; summed generation 3302.459 s, submission/fallback 4395.356 s; wall 481.155 s
Frame 0: catalog splatting finished in 481.2 s; writing image...
Rendered 598264924 images from 343363141 catalog stars to output/imgs/2mass_galactic_center_schwarzschild_R100_4k_1e13_beta4.5_16_4-full-gpu.png (ok; imported lens map)
PSF splats: cached 598264924, cached wing-clipped 0, direct fallbacks 0, discarded below min-Y 0
HIP PSF: 598264924 events in 36526 batches (36526 timed); upload 1.504570 s, kernel 476.906037 s, download 0.021966 s
./build/Release/schwarzschild_sky_hip --verbose --psf-relative-tail 1e-8 0 7733.61s user 61.18s system 1462% cpu 8:53.06 total
benchmarks/hip_psf_2026-09-07.md
+204
View File
@@ -0,0 +1,204 @@
# HIP PSF 有界排查与优化（2026-09-07）
后续用户完成了全天空 CPU/HIP 配对运行，并确认 PNG 比特级一致；见
[完整 4K 记录](2mass_galactic_center_cpu_hip_2026-09-07.md)，其中记录了本次优化提交的实际 hash。
本页保留提交前有界排查的原始上下文。
## 范围与结论
本次没有运行用户的完整 343,363,141 星 / 598,308,919 星像任务。
所有 renderer/benchmark 都单进程顺序执行，前一个退出后才启动下一个；单次微基准
45 秒超时，缩小 renderer 60 秒超时。没有并发性能测试，没有删除原 render 输出。
源码基线为 `265d7b95d5b916ec4768d174d73ad963a43f0088`，优化版是此基线上的未提交改动。
最初保留了用户现存 HIP binary；另从该 Git commit 的干净源码重建旧版，以复核对照。
主机 HIP 7.2.53210-9999，AMD Radeon RX 9070 / gfx1201；性能测试设置
`OMP_NUM_THREADS=16`，catalog loaders 为 4。Release / `-O2`；HIP 显式
`--offload-arch=gfx1201`。完整元数据见 [metadata.log](hip_psf_2026-09-07/metadata.log)。
已确认的实现问题与处理：
- 旧 HIP 分支将 catalog 查询、逆映射、频移与黑体积分串行化。现在按 triangle 使用
OpenMP producer，每 worker 一个 16,384-event 私有块，GPU cache/HDR 只分配一次。
- 旧 kernel 一线程串行遍历一个 PSF，double atomic 编译为 64-bit CAS 重试循环。
现在每个 event 用 32 线程访问相邻像素；保持四点 phase cache、圆形 support 和
double RGB 累积。分布实验支持这一改变，不把全部 kernel 时间归因为 atomic 竞争。
- 旧 host 缓冲是可立即复用的普通 malloc 数组；现在 sink 同步复制进自己拥有的两个
pinned slot，重用 host/device slot 前等待上一 kernel 完成，生命周期不依赖 pageable
`hipMemcpyAsync` 的运行时暂存行为。GPU copy/kernel 仍在同一 stream 内顺序执行。
- 旧计时只覆盖前 64 批。现在两个 slot 的计时 event 循环使用，统计全部批次，并约每
5 秒在提交时打印进度。完成数是已回收计时的完成数，最多有两个未回收 slot。
- chunk 提交与整个 CPU direct-fallback 的 download/evaluate/upload 持有同一共享锁。
没有 per-star 锁，没有共享可变 CPU scratch，没有改变 backend 错误时立即失败的行为。
未使用分桶规约、float HDR、近似黑体积分或改变 min-Y/tail 参数。
## 受控 kernel 实验
使用 `tests/benchmark_hip_psf.hip` 生成固定 LCG seed 12345 的相同事件。frame 3840×2160，
FWHM 2.7，beta 4.5，tail 1e-8，cache radius 47；每个事件 56 B。
每个进程先计算串行 CPU double HDR reference，再执行一个 HIP replay。CPU reference
耗时不是 production 16-thread CPU render 的性能基线。
| 事件数 | 中心分布范围 | 旧 kernel | 32-thread kernel 原型 | 比值 |
| ---: | ---: | ---: | ---: | ---: |
| 4,096 | 16 px | 22.490 ms | 5.702 ms | 3.94× |
| 16,384 | 1 px | 137.888 ms | 23.976 ms | 5.75× |
| 16,384 | 3840 px（y 限于2160） | 55.630 ms | 21.492 ms | 2.59× |
原型阶段仅改变 kernel 的线程分工，保留旧上传和计时实现，便于隔离效果。
原始输出：[旧密集](hip_psf_2026-09-07/baseline_dense.log)、
[新密集](hip_psf_2026-09-07/wave_dense.log)、
[旧热点](hip_psf_2026-09-07/baseline_hotspot.log)、
[新热点](hip_psf_2026-09-07/wave_hotspot.log)、
[旧分散](hip_psf_2026-09-07/baseline_spread.log)、
[新分散](hip_psf_2026-09-07/wave_spread.log)。
整合双缓冲后的 65,536-event/16px replay：4/4 批均计时，kernel 78.268 ms，replay wall
83.762 ms；CPU reference 842.402 ms。double HDR max_abs `1.7053025658242404e-13`，
max_rel `3.3472443502765663e-15`，总 RGB flux 相对差 `-1.3316774058238102e-17`。
[原始输出](hip_psf_2026-09-07/final_multibatch.log)。
## 原 4K lens map 的真实 catalog 子集
保留原文件 `output/lens/schwarzschild_galactic_center_R100_45deg_16_4_4k_.grlens`：
3840×2160、45°、65,817 vertices、130,882 triangles。未重新 tracing/refine。
使用 tile 索引的主对照输入只有三个 tile：RA252/Dec048、RA262/Dec060、RA082/Dec090。
每个 tile 均匀抽取最多 8192 行，共 21,087 颗星，产生 156,837 个星像。
准备脚本和源/子集 SHA256：[prepare_inputs.py](hip_psf_2026-09-07/prepare_inputs.py)、
[input_manifest.log](hip_psf_2026-09-07/input_manifest.log)。
其余 tile 故意缺失，因此 prefetch 的 64,437 unavailable 是缩小输入条件，不是生产数据错误。
| 实现 | catalog/splat（日志四舍五入） | GPU kernel 合计 | 完整进程 |
| --- | ---: | ---: | ---: |
| 原有 HIP binary | 1.9 s | 1.469529 s | 2.971 s |
| 干净旧源码重建 HIP | 1.9 s | 1.496493 s | 3.016 s |
| 优化 HIP | 0.4 s | 0.168641 s | 1.487 s |
| CPU 16 workers | 0.6 s | — | 1.674 s |
所有实现均 cached=156837、wing-clipped=0、direct=0、discarded=0。
优化版 GPU 16 批全部计时；producer wall=0.326 s，worker generation 总和=2.053 s，
submission/fallback 总和=1.157 s。worker wall 总和包含并发时间，不能加到 GPU 时间上。
原始日志：[旧 HIP](hip_psf_2026-09-07/baseline_multi.log)、
[优化 HIP](hip_psf_2026-09-07/final_multi.log)、[CPU](hip_psf_2026-09-07/cpu_multi.log)。
干净源码旧版复核另见 [clean_baseline_multi.log](hip_psf_2026-09-07/clean_baseline_multi.log)。
旧 HIP 与优化 HIP 的 kernel 时间相差约 8.7 倍、完整进程约 2 倍；这些是小规模单次测量，
不是完整全天空的加速比。相较 CPU，本子集的完整进程收益较小，cache 初始化和 4K 输出
占比较高。没有依此预测用户完整任务的完成时间。
辅助诊断保留了同一 lens map 的单 CSV 模式：256 星及 8192 星。
8192 星时旧/新 splat 为 4.4/0.6 秒，完整进程 5.489/1.706 秒。
但单 CSV 会放大逐 triangle 扫描成本，不能用它代表生产 tile 索引的 CPU 占比。
同 8192 星改用 tile 索引后，旧 kernel=0.171626 秒，新 kernel=0.028882 秒。
日志：[256](hip_psf_2026-09-07/baseline_lens256.log)、
[CSV旧](hip_psf_2026-09-07/baseline_lens8192.log)、
[CSV新](hip_psf_2026-09-07/final_lens8192.log)、
[索引旧](hip_psf_2026-09-07/baseline_indexed.log)、
[索引新](hip_psf_2026-09-07/final_indexed.log)。
## 验证
- `make -j1 test`：通过；固定 HDR、camera、geodesic、frame、Schwarzschild、observer track、
catalog prefetch、CLI 和多线程 movie/lens map 回归均通过。
[完整日志](hip_psf_2026-09-07/cpu_regression.log)。
- `test_hip_psf`：213 事件、71 次 submit，全部计时；覆盖 host 数组立即覆写、slot 复用、
画面边缘、wing clipping 和中途 CPU HDR 修改后继续 GPU。
[初次日志](hip_psf_2026-09-07/streaming_test.log)、
[最终重建后日志](hip_psf_2026-09-07/streaming_test_final.log)。
- `tests/test_hip_renderer.py`：96×64、4 workers，先 CPU 后 HIP；验证 cached=2/direct=1
的混合场和 discarded=3 的 min-Y 场。统计一致、FITS float32 HDR payload 完全一致。
[完整日志](hip_psf_2026-09-07/renderer_regression.log)。
- 4K 子集的旧/新 HIP 以及 CPU/新 HIP：24,883,200 个 FITS float32 样本完全一致；PNG
也一致。double 精度由上述 sink/replay tests 单独验证，未将 FITS 当作 double 文件。
[比较日志](hip_psf_2026-09-07/hdr_comparison.log)。
- 在单个子进程设置 `ROCR_VISIBLE_DEVICES=''`，32×24 的显式 HIP render 报设备不可用、
非零退出且没有写图。[原始输出](hip_psf_2026-09-07/fail_fast.log)。
- 初次 mixed fixture 使用 exposure=2 未越过 support 阈值，测试断言正确拒绝了无 fallback
的输入；依据源码改为 exposure=1000 后才通过。这是测试输入校正，非 production 修复。
[初次输出](hip_psf_2026-09-07/fixture_calibration_failed.log)。
## 可复制命令
以下使用 zsh 默认 `time`，不设置 GNU 风格 `TIMEFMT`。命令必须顺序执行；任何失败或
超时后先确认前一个进程已退出。原始输出保存在上述链接，未只保留汇总值。
```zsh
# 从本仓库根目录执行；准备固定样本，输出源与子集哈希。
python3 benchmarks/hip_psf_2026-09-07/prepare_inputs.py
# 独立重建旧源码，不改变当前 checkout。
mkdir -p /tmp/gr-hip-baseline-src
git archive 265d7b95d5b916ec4768d174d73ad963a43f0088 Makefile src mk | tar -x -C /tmp/gr-hip-baseline-src
make -j4 -C /tmp/gr-hip-baseline-src PSF_BACKEND=hip SPACETIME=schwarzschild ENABLE_HDR=1 \
HIP_CXXFLAGS='--offload-arch=gfx1201 -std=c++17 -O2 -Wall -Wextra -Wpedantic' backend
# 当前实现与 CPU reference。
make -j4 PSF_BACKEND=hip SPACETIME=schwarzschild ENABLE_HDR=1 \
HIP_CXXFLAGS='--offload-arch=gfx1201 -std=c++17 -O2 -Wall -Wextra -Wpedantic' backend hip-psf-test hip-psf-bench
make -j4 PSF_BACKEND=cpu SPACETIME=schwarzschild ENABLE_HDR=1 backend
render_args=(
--all-sky-catalog /tmp/gr-hip-catalog-multi --catalog-load-workers 4
--psf-relative-tail 1e-8 --psf-min-y 0 --max-cache-psf-flux 1e8
--exposure 1e13 --psf-fwhm-pixels 2.7 --psf-moffat-beta 4.5 --hdr-output
--lens-map-input output/lens/schwarzschild_galactic_center_R100_45deg_16_4_4k_.grlens
--verbose
)
time OMP_NUM_THREADS=16 timeout --kill-after=5s 60s /tmp/gr-hip-baseline-src/build/Release/schwarzschild_sky_hip \
"${render_args[@]}" --output /tmp/gr-hip-clean-baseline-multi.png
time OMP_NUM_THREADS=16 timeout --kill-after=5s 60s ./build/Release/schwarzschild_sky_hip \
"${render_args[@]}" --output /tmp/gr-hip-final-multi.png
time OMP_NUM_THREADS=16 timeout --kill-after=5s 60s ./build/Release/schwarzschild_sky \
"${render_args[@]}" --output /tmp/gr-cpu-multi.png
# 原有 binary 的初始记录使用 /tmp/gr-schwarzschild-hip-baseline，其他参数完全相同。
# 单 tile 索引测试仅将 root 换为 /tmp/gr-hip-catalog-subset。
# 单 CSV 诊断改用 --catalog benchmarks/hip_psf_2026-09-07/catalog_8192.csv
# （或 catalog_256.csv），不用 --all-sky-catalog，输出到单独的 /tmp PNG。
OMP_NUM_THREADS=16 timeout --kill-after=5s 45s ./build/Release/benchmark_hip_psf 65536 16
OMP_NUM_THREADS=16 timeout --kill-after=5s 45s ./build/Release/test_hip_psf
make -j1 test
# 小尺寸 renderer 回归；先分别构建带 HDR 的 Minkowski CPU/HIP。
make PSF_BACKEND=cpu SPACETIME=minkowski ENABLE_HDR=1 backend
make PSF_BACKEND=hip SPACETIME=minkowski ENABLE_HDR=1 \
HIP_CXXFLAGS='--offload-arch=gfx1201 -std=c++17 -O2 -Wall -Wextra -Wpedantic' backend
python3 tests/test_hip_renderer.py build/Release/minkowski_sky build/Release/minkowski_sky_hip
python3 scripts/fits_floatdiff.py /tmp/gr-cpu-multi_HDR.fits /tmp/gr-hip-final-multi_HDR.fits
```
旧/新原型微基准可使用同一 `tests/benchmark_hip_psf.hip` 与基线 sink 源码编译：
```zsh
cp /tmp/gr-hip-baseline-src/src/hip_psf.hip /tmp/gr-hip-psf-baseline.hip
hipcc --offload-arch=gfx1201 -std=c++17 -O2 -fopenmp -Isrc -x hip \
tests/benchmark_hip_psf.hip /tmp/gr-hip-psf-baseline.hip -x none \
build/Release/obj/schwarzschild/hdr_sink1_hip/src/optics.o \
-lm -lpng -lcfitsio -o /tmp/hip_psf_bench_baseline
python3 - <<'PY'
from pathlib import Path
s = Path('/tmp/gr-hip-psf-baseline.hip').read_text()
s = s.replace('const size_t index = (size_t)blockIdx.x * blockDim.x + threadIdx.x;',
'const size_t thread = (size_t)blockIdx.x * blockDim.x + threadIdx.x;\n'
' const size_t index = thread / 32;\n const int lane = thread % 32;')
s = s.replace('int offset_x = first; offset_x <= last; ++offset_x',
'int offset_x = first + lane; offset_x <= last; offset_x += 32')
s = s.replace('(event_count + threads - 1) / threads',
'(event_count * 32 + threads - 1) / threads')
Path('/tmp/gr-hip-psf-wave.hip').write_text(s)
PY
hipcc --offload-arch=gfx1201 -std=c++17 -O2 -fopenmp -Isrc -x hip \
tests/benchmark_hip_psf.hip /tmp/gr-hip-psf-wave.hip -x none \
build/Release/obj/schwarzschild/hdr_sink1_hip/src/optics.o \
-lm -lpng -lcfitsio -o /tmp/hip_psf_bench_wave
OMP_NUM_THREADS=16 timeout --kill-after=5s 45s /tmp/hip_psf_bench_baseline 4096 16
OMP_NUM_THREADS=16 timeout --kill-after=5s 45s /tmp/hip_psf_bench_wave 4096 16
OMP_NUM_THREADS=16 timeout --kill-after=5s 45s /tmp/hip_psf_bench_baseline 16384 1
OMP_NUM_THREADS=16 timeout --kill-after=5s 45s /tmp/hip_psf_bench_wave 16384 1
OMP_NUM_THREADS=16 timeout --kill-after=5s 45s /tmp/hip_psf_bench_baseline 16384 3840
OMP_NUM_THREADS=16 timeout --kill-after=5s 45s /tmp/hip_psf_bench_wave 16384 3840
```
benchmarks/hip_psf_2026-09-07/baseline_dense.log
+3
View File
@@ -0,0 +1,3 @@
PSF cache ready: 64x64 phases, radius 47 px, relative tail 1e-08, tail abs 1e-06, boundary 1e-07, build 0.661 s
device=AMD Radeon RX 9070 arch=gfx1201 events=4096 spread=16 frame=3840x2160 event_bytes=56 CPU_reference=0.051952 s
HIP create=0.027773 replay_wall=0.027614 upload=0.000021 kernel=0.022490 download=0.004796 batches=1 timed=1 max_abs=4.7073456244106637e-14 max_rel=1.1345618109894268e-14 flux_rel=1.138425198954238e-16
benchmarks/hip_psf_2026-09-07/baseline_hotspot.log
+3
View File
@@ -0,0 +1,3 @@
PSF cache ready: 64x64 phases, radius 47 px, relative tail 1e-08, tail abs 1e-06, boundary 1e-07, build 0.681 s
device=AMD Radeon RX 9070 arch=gfx1201 events=16384 spread=1 frame=3840x2160 event_bytes=56 CPU_reference=0.210148 s
HIP create=0.025685 replay_wall=0.143944 upload=0.000049 kernel=0.137888 download=0.005726 batches=1 timed=1 max_abs=2.6147972675971687e-12 max_rel=1.6347773919004679e-14 flux_rel=-2.3341330284619701e-16
benchmarks/hip_psf_2026-09-07/baseline_indexed.log
+10
View File
@@ -0,0 +1,10 @@
PSF cache ready: 64x64 phases, radius 47 px, relative tail 1e-08, tail abs 1e-06, boundary 1e-07, build 0.836 s
Frame 0: finding and prefetching catalog tiles...
Catalog prefetch: 64440 requested, 1 newly loaded (8192 stars), 64439 unavailable in 0.033 s; 4 loader workers
Frame 0: catalog prefetch finished (64440 candidate tiles).
Frame 0: splatting 130882 lens triangles...
Frame 0: catalog splatting finished in 0.3 s; writing image...
Rendered 12765 images from 8192 catalog stars to /tmp/gr-hip-baseline-indexed.png (ok; imported lens map)
PSF splats: cached 12765, cached wing-clipped 0, direct fallbacks 0, discarded below min-Y 0
HIP PSF: 12765 events in 1 batches (1 timed); upload 0.000042 s, kernel 0.171626 s, download 0.025845 s
OMP_NUM_THREADS=16 timeout --kill-after=5s 60s --all-sky-catalog 4 1e-8 11.23s user 0.26s system 721% cpu 1.592 total
benchmarks/hip_psf_2026-09-07/baseline_lens256.log
+9
View File
@@ -0,0 +1,9 @@
PSF cache ready: 64x64 phases, radius 47 px, relative tail 1e-08, tail abs 1e-06, boundary 1e-07, build 0.709 s
Frame 0: finding and prefetching catalog tiles...
Frame 0: catalog prefetch finished (0 candidate tiles).
Frame 0: splatting 130882 lens triangles...
Frame 0: catalog splatting finished in 0.3 s; writing image...
Rendered 411 images from 256 catalog stars to /tmp/gr-hip-baseline-256.png (ok; imported lens map)
PSF splats: cached 411, cached wing-clipped 0, direct fallbacks 0, discarded below min-Y 0
HIP PSF: 411 events in 1 batches (1 timed); upload 0.000014 s, kernel 0.122597 s, download 0.021485 s
OMP_NUM_THREADS=16 timeout --kill-after=5s 60s --catalog --psf-relative-tai 10.69s user 0.17s system 774% cpu 1.403 total
benchmarks/hip_psf_2026-09-07/baseline_lens8192.log
+9
View File
@@ -0,0 +1,9 @@
PSF cache ready: 64x64 phases, radius 47 px, relative tail 1e-08, tail abs 1e-06, boundary 1e-07, build 0.694 s
Frame 0: finding and prefetching catalog tiles...
Frame 0: catalog prefetch finished (0 candidate tiles).
Frame 0: splatting 130882 lens triangles...
Frame 0: catalog splatting finished in 4.4 s; writing image...
Rendered 12765 images from 8192 catalog stars to /tmp/gr-hip-baseline-8192.png (ok; imported lens map)
PSF splats: cached 12765, cached wing-clipped 0, direct fallbacks 0, discarded below min-Y 0
HIP PSF: 12765 events in 1 batches (1 timed); upload 0.000041 s, kernel 0.171573 s, download 0.025041 s
OMP_NUM_THREADS=16 timeout --kill-after=5s 60s --catalog --psf-relative-tai 14.92s user 0.17s system 274% cpu 5.489 total
benchmarks/hip_psf_2026-09-07/baseline_multi.log
+10
View File
@@ -0,0 +1,10 @@
PSF cache ready: 64x64 phases, radius 47 px, relative tail 1e-08, tail abs 1e-06, boundary 1e-07, build 0.674 s
Frame 0: finding and prefetching catalog tiles...
Catalog prefetch: 64440 requested, 3 newly loaded (21087 stars), 64437 unavailable in 0.033 s; 4 loader workers
Frame 0: catalog prefetch finished (64440 candidate tiles).
Frame 0: splatting 130882 lens triangles...
Frame 0: catalog splatting finished in 1.9 s; writing image...
Rendered 156837 images from 21087 catalog stars to /tmp/gr-hip-baseline-multi.png (ok; imported lens map)
PSF splats: cached 156837, cached wing-clipped 0, direct fallbacks 0, discarded below min-Y 0
HIP PSF: 156837 events in 10 batches (10 timed); upload 0.000312 s, kernel 1.469529 s, download 0.025304 s
OMP_NUM_THREADS=16 timeout --kill-after=5s 60s --all-sky-catalog 4 1e-8 12.51s user 0.24s system 429% cpu 2.971 total
benchmarks/hip_psf_2026-09-07/baseline_spread.log
+3
View File
@@ -0,0 +1,3 @@
PSF cache ready: 64x64 phases, radius 47 px, relative tail 1e-08, tail abs 1e-06, boundary 1e-07, build 0.677 s
device=AMD Radeon RX 9070 arch=gfx1201 events=16384 spread=3840 frame=3840x2160 event_bytes=56 CPU_reference=0.419966 s
HIP create=0.023973 replay_wall=0.060952 upload=0.000049 kernel=0.055630 download=0.004973 batches=1 timed=1 max_abs=5.5511151231257827e-17 max_rel=1.0443479008602178e-15 flux_rel=-5.9614463770241093e-20
benchmarks/hip_psf_2026-09-07/build_baseline_clean.log
+16
View File
@@ -0,0 +1,16 @@
make: Entering directory '/tmp/gr-hip-baseline-src'
cc -DENABLE_PNG -DFRAME_PSF_EVENT_SINK=1 -DPSF_BACKEND_HIP -DSPACETIME_SCHWARZSCHILD -DENABLE_HDR_OUTPUT -std=c11 -march=native -pipe -Wall -Wextra -Wpedantic -O2 -DNDEBUG -fopenmp -Isrc -MMD -MP -c src/catalog.c -o build/Release/obj/schwarzschild/hdr_sink1_hip/src/catalog.o
cc -DENABLE_PNG -DFRAME_PSF_EVENT_SINK=1 -DPSF_BACKEND_HIP -DSPACETIME_SCHWARZSCHILD -DENABLE_HDR_OUTPUT -std=c11 -march=native -pipe -Wall -Wextra -Wpedantic -O2 -DNDEBUG -fopenmp -Isrc -MMD -MP -c src/frame.c -o build/Release/obj/schwarzschild/hdr_sink1_hip/src/frame.o
cc -DENABLE_PNG -DFRAME_PSF_EVENT_SINK=1 -DPSF_BACKEND_HIP -DSPACETIME_SCHWARZSCHILD -DENABLE_HDR_OUTPUT -std=c11 -march=native -pipe -Wall -Wextra -Wpedantic -O2 -DNDEBUG -fopenmp -Isrc -MMD -MP -c src/geodesic.c -o build/Release/obj/schwarzschild/hdr_sink1_hip/src/geodesic.o
cc -DENABLE_PNG -DFRAME_PSF_EVENT_SINK=1 -DPSF_BACKEND_HIP -DSPACETIME_SCHWARZSCHILD -DENABLE_HDR_OUTPUT -std=c11 -march=native -pipe -Wall -Wextra -Wpedantic -O2 -DNDEBUG -fopenmp -Isrc -MMD -MP -c src/lens_map.c -o build/Release/obj/schwarzschild/hdr_sink1_hip/src/lens_map.o
cc -DENABLE_PNG -DFRAME_PSF_EVENT_SINK=1 -DPSF_BACKEND_HIP -DSPACETIME_SCHWARZSCHILD -DENABLE_HDR_OUTPUT -std=c11 -march=native -pipe -Wall -Wextra -Wpedantic -O2 -DNDEBUG -fopenmp -Isrc -MMD -MP -c src/movie.c -o build/Release/obj/schwarzschild/hdr_sink1_hip/src/movie.o
cc -DENABLE_PNG -DFRAME_PSF_EVENT_SINK=1 -DPSF_BACKEND_HIP -DSPACETIME_SCHWARZSCHILD -DENABLE_HDR_OUTPUT -std=c11 -march=native -pipe -Wall -Wextra -Wpedantic -O2 -DNDEBUG -fopenmp -Isrc -MMD -MP -c src/observer.c -o build/Release/obj/schwarzschild/hdr_sink1_hip/src/observer.o
cc -DENABLE_PNG -DFRAME_PSF_EVENT_SINK=1 -DPSF_BACKEND_HIP -DSPACETIME_SCHWARZSCHILD -DENABLE_HDR_OUTPUT -std=c11 -march=native -pipe -Wall -Wextra -Wpedantic -O2 -DNDEBUG -fopenmp -Isrc -MMD -MP -c src/observer_track.c -o build/Release/obj/schwarzschild/hdr_sink1_hip/src/observer_track.o
cc -DENABLE_PNG -DFRAME_PSF_EVENT_SINK=1 -DPSF_BACKEND_HIP -DSPACETIME_SCHWARZSCHILD -DENABLE_HDR_OUTPUT -std=c11 -march=native -pipe -Wall -Wextra -Wpedantic -O2 -DNDEBUG -fopenmp -Isrc -MMD -MP -c src/optics.c -o build/Release/obj/schwarzschild/hdr_sink1_hip/src/optics.o
cc -DENABLE_PNG -DFRAME_PSF_EVENT_SINK=1 -DPSF_BACKEND_HIP -DSPACETIME_SCHWARZSCHILD -DENABLE_HDR_OUTPUT -std=c11 -march=native -pipe -Wall -Wextra -Wpedantic -O2 -DNDEBUG -fopenmp -Isrc -MMD -MP -c src/ray.c -o build/Release/obj/schwarzschild/hdr_sink1_hip/src/ray.o
cc -DENABLE_PNG -DFRAME_PSF_EVENT_SINK=1 -DPSF_BACKEND_HIP -DSPACETIME_SCHWARZSCHILD -DENABLE_HDR_OUTPUT -std=c11 -march=native -pipe -Wall -Wextra -Wpedantic -O2 -DNDEBUG -fopenmp -Isrc -MMD -MP -c src/spacetime_common.c -o build/Release/obj/schwarzschild/hdr_sink1_hip/src/spacetime_common.o
cc -DENABLE_PNG -DFRAME_PSF_EVENT_SINK=1 -DPSF_BACKEND_HIP -DSPACETIME_SCHWARZSCHILD -DENABLE_HDR_OUTPUT -std=c11 -march=native -pipe -Wall -Wextra -Wpedantic -O2 -DNDEBUG -fopenmp -Isrc -MMD -MP -c src/spacetime_schwarzschild.c -o build/Release/obj/schwarzschild/hdr_sink1_hip/src/spacetime_schwarzschild.o
cc -DENABLE_PNG -DFRAME_PSF_EVENT_SINK=1 -DPSF_BACKEND_HIP -DSPACETIME_SCHWARZSCHILD -DENABLE_HDR_OUTPUT -std=c11 -march=native -pipe -Wall -Wextra -Wpedantic -O2 -DNDEBUG -fopenmp -Isrc -MMD -MP -c src/main.c -o build/Release/obj/schwarzschild/hdr_sink1_hip/src/main.o
hipcc --offload-arch=gfx1201 -std=c++17 -O2 -Wall -Wextra -Wpedantic -DFRAME_PSF_EVENT_SINK=1 -DPSF_BACKEND_HIP -Isrc -c src/hip_psf.hip -o build/Release/obj/schwarzschild/hdr_sink1_hip/src/hip_psf.o
hipcc -O2 -DNDEBUG -fopenmp build/Release/obj/schwarzschild/hdr_sink1_hip/src/catalog.o build/Release/obj/schwarzschild/hdr_sink1_hip/src/frame.o build/Release/obj/schwarzschild/hdr_sink1_hip/src/geodesic.o build/Release/obj/schwarzschild/hdr_sink1_hip/src/lens_map.o build/Release/obj/schwarzschild/hdr_sink1_hip/src/movie.o build/Release/obj/schwarzschild/hdr_sink1_hip/src/observer.o build/Release/obj/schwarzschild/hdr_sink1_hip/src/observer_track.o build/Release/obj/schwarzschild/hdr_sink1_hip/src/optics.o build/Release/obj/schwarzschild/hdr_sink1_hip/src/ray.o build/Release/obj/schwarzschild/hdr_sink1_hip/src/spacetime_common.o build/Release/obj/schwarzschild/hdr_sink1_hip/src/spacetime_schwarzschild.o build/Release/obj/schwarzschild/hdr_sink1_hip/src/main.o build/Release/obj/schwarzschild/hdr_sink1_hip/src/hip_psf.o -lm -lpng -lcfitsio -o build/Release/schwarzschild_sky_hip
make: Leaving directory '/tmp/gr-hip-baseline-src'
benchmarks/hip_psf_2026-09-07/build_cpu.log
+2
View File
@@ -0,0 +1,2 @@
cc -DENABLE_PNG -DFRAME_PSF_EVENT_SINK=1 -DSPACETIME_SCHWARZSCHILD -DENABLE_HDR_OUTPUT -std=c11 -march=native -pipe -Wall -Wextra -Wpedantic -O2 -DNDEBUG -fopenmp -Isrc -MMD -MP -c src/frame.c -o build/Release/obj/schwarzschild/hdr_sink1_cpu/src/frame.o
cc -O2 -DNDEBUG -fopenmp build/Release/obj/schwarzschild/hdr_sink1_cpu/src/catalog.o build/Release/obj/schwarzschild/hdr_sink1_cpu/src/frame.o build/Release/obj/schwarzschild/hdr_sink1_cpu/src/geodesic.o build/Release/obj/schwarzschild/hdr_sink1_cpu/src/lens_map.o build/Release/obj/schwarzschild/hdr_sink1_cpu/src/movie.o build/Release/obj/schwarzschild/hdr_sink1_cpu/src/observer.o build/Release/obj/schwarzschild/hdr_sink1_cpu/src/observer_track.o build/Release/obj/schwarzschild/hdr_sink1_cpu/src/optics.o build/Release/obj/schwarzschild/hdr_sink1_cpu/src/ray.o build/Release/obj/schwarzschild/hdr_sink1_cpu/src/spacetime_common.o build/Release/obj/schwarzschild/hdr_sink1_cpu/src/spacetime_schwarzschild.o build/Release/obj/schwarzschild/hdr_sink1_cpu/src/main.o -lm -lpng -lcfitsio -o build/Release/schwarzschild_sky
benchmarks/hip_psf_2026-09-07/build_hip.log
+3
View File
@@ -0,0 +1,3 @@
hipcc -O2 -DNDEBUG -fopenmp build/Release/obj/schwarzschild/hdr_sink1_hip/src/catalog.o build/Release/obj/schwarzschild/hdr_sink1_hip/src/frame.o build/Release/obj/schwarzschild/hdr_sink1_hip/src/geodesic.o build/Release/obj/schwarzschild/hdr_sink1_hip/src/lens_map.o build/Release/obj/schwarzschild/hdr_sink1_hip/src/movie.o build/Release/obj/schwarzschild/hdr_sink1_hip/src/observer.o build/Release/obj/schwarzschild/hdr_sink1_hip/src/observer_track.o build/Release/obj/schwarzschild/hdr_sink1_hip/src/optics.o build/Release/obj/schwarzschild/hdr_sink1_hip/src/ray.o build/Release/obj/schwarzschild/hdr_sink1_hip/src/spacetime_common.o build/Release/obj/schwarzschild/hdr_sink1_hip/src/spacetime_schwarzschild.o build/Release/obj/schwarzschild/hdr_sink1_hip/src/main.o build/Release/obj/schwarzschild/hdr_sink1_hip/src/hip_psf.o -lm -lpng -lcfitsio -o build/Release/schwarzschild_sky_hip
make: Nothing to be done for 'hip-psf-test'.
hipcc --offload-arch=gfx1201 -std=c++17 -O2 -Wall -Wextra -Wpedantic -fopenmp -Isrc -x hip tests/benchmark_hip_psf.hip -x none build/Release/obj/schwarzschild/hdr_sink1_hip/src/hip_psf.o build/Release/obj/schwarzschild/hdr_sink1_hip/src/optics.o -lm -lpng -lcfitsio -o build/Release/benchmark_hip_psf
benchmarks/hip_psf_2026-09-07/build_hip_final.log
+4
View File
@@ -0,0 +1,4 @@
hipcc --offload-arch=gfx1201 -std=c++17 -O2 -Wall -Wextra -Wpedantic -DFRAME_PSF_EVENT_SINK=1 -DPSF_BACKEND_HIP -Isrc -c src/hip_psf.hip -o build/Release/obj/schwarzschild/hdr_sink1_hip/src/hip_psf.o
hipcc -O2 -DNDEBUG -fopenmp build/Release/obj/schwarzschild/hdr_sink1_hip/src/catalog.o build/Release/obj/schwarzschild/hdr_sink1_hip/src/frame.o build/Release/obj/schwarzschild/hdr_sink1_hip/src/geodesic.o build/Release/obj/schwarzschild/hdr_sink1_hip/src/lens_map.o build/Release/obj/schwarzschild/hdr_sink1_hip/src/movie.o build/Release/obj/schwarzschild/hdr_sink1_hip/src/observer.o build/Release/obj/schwarzschild/hdr_sink1_hip/src/observer_track.o build/Release/obj/schwarzschild/hdr_sink1_hip/src/optics.o build/Release/obj/schwarzschild/hdr_sink1_hip/src/ray.o build/Release/obj/schwarzschild/hdr_sink1_hip/src/spacetime_common.o build/Release/obj/schwarzschild/hdr_sink1_hip/src/spacetime_schwarzschild.o build/Release/obj/schwarzschild/hdr_sink1_hip/src/main.o build/Release/obj/schwarzschild/hdr_sink1_hip/src/hip_psf.o -lm -lpng -lcfitsio -o build/Release/schwarzschild_sky_hip
hipcc --offload-arch=gfx1201 -std=c++17 -O2 -Wall -Wextra -Wpedantic -fopenmp -Isrc -x hip tests/test_hip_psf.hip -x none build/Release/obj/schwarzschild/hdr_sink1_hip/src/hip_psf.o build/Release/obj/schwarzschild/hdr_sink1_hip/src/optics.o -lm -lpng -lcfitsio -o build/Release/test_hip_psf
hipcc --offload-arch=gfx1201 -std=c++17 -O2 -Wall -Wextra -Wpedantic -fopenmp -Isrc -x hip tests/benchmark_hip_psf.hip -x none build/Release/obj/schwarzschild/hdr_sink1_hip/src/hip_psf.o build/Release/obj/schwarzschild/hdr_sink1_hip/src/optics.o -lm -lpng -lcfitsio -o build/Release/benchmark_hip_psf
benchmarks/hip_psf_2026-09-07/catalog_256.csv
+257
View File
@@ -0,0 +1,257 @@
ra_deg,dec_deg,temperature_K,amplitude_sr
252.995803,-41.503754,2122.6076,8.520470492e-21
252.993551,-41.543816,2114.6385,9.967579356e-21
252.918388,-41.510990,2208.128,1.446557073e-19
252.900742,-41.506279,2279.2582,1.973490936e-18
252.884436,-41.488571,2138.1302,1.410128710e-20
252.986092,-41.597309,2129.5884,3.136569365e-20
252.984711,-41.606659,2114.2229,9.627036994e-21
252.924427,-41.399944,3134.4748,2.266711014e-21
252.938142,-41.386372,5574.2493,7.758286100e-21
252.958086,-41.625877,1617.2207,3.144977882e-19
252.970634,-41.365410,2202.8994,1.648245643e-20
252.810573,-41.475826,3941.0552,1.524175586e-21
252.854571,-41.396477,2408.7178,2.646847276e-20
252.974001,-41.655529,2217.0089,4.616359171e-20
252.783806,-41.501965,2321.7548,7.149940580e-21
252.782782,-41.459801,2703.1671,4.949843960e-20
252.880635,-41.648254,2158.5784,8.959667114e-21
252.832964,-41.372025,2223.318,6.324985653e-20
252.781331,-41.583836,4289.9357,3.687496634e-21
252.774616,-41.415287,2503.0345,3.528731590e-20
252.781981,-41.604404,4481.0767,4.210660307e-21
252.835709,-41.344860,2471.4445,7.859657597e-20
252.755037,-41.586941,2520.0282,9.071531178e-21
252.829319,-41.663242,2092.7219,9.439286836e-21
252.729244,-41.560253,3870.9986,2.202578638e-21
252.792426,-41.349979,2347.9017,6.933492283e-21
252.715858,-41.556286,3197.6305,2.002903960e-21
252.745823,-41.618793,2093.3838,1.798894237e-20
252.869507,-41.293640,2565.0983,7.621517442e-21
252.968856,-41.268589,2990.034,2.505403621e-21
252.925736,-41.270466,3338.3343,2.122223459e-21
252.680916,-41.525299,1724.5512,7.687633387e-18
252.701932,-41.401005,2431.4439,5.080775898e-20
252.829047,-41.712399,2045.1192,5.651479174e-20
252.678461,-41.426064,3910.3355,1.929630882e-21
252.944337,-41.247566,2738.4573,2.298467128e-20
252.821892,-41.277275,2590.152,3.408118915e-21
252.662337,-41.573349,2041.2244,1.499212709e-20
252.651190,-41.553684,3393.4186,2.240876206e-21
252.717832,-41.332279,2490.4613,6.309850262e-20
252.859629,-41.752144,1851.5385,1.237866868e-19
252.710481,-41.329098,3931.3812,2.005228197e-21
252.712726,-41.321545,3857.1544,5.275918149e-21
252.907975,-41.225502,2783.4349,7.275926125e-21
252.617463,-41.497688,2452.5593,1.840842057e-20
252.697290,-41.680248,2279.1812,1.329110424e-20
252.701091,-41.312347,2287.752,8.783238970e-21
252.641947,-41.624428,2447.7312,1.460562478e-20
252.611005,-41.565239,2414.9482,7.249220492e-21
252.921863,-41.204163,2880.0791,4.126485206e-21
252.639833,-41.358757,3812.7671,1.743956802e-21
252.594286,-41.548817,3760.853,2.953969609e-21
252.630135,-41.641548,2306.0382,3.363985170e-19
252.825565,-41.215195,2246.365,7.392177336e-20
252.620751,-41.361652,2520.4538,1.047956906e-20
252.802403,-41.217541,4186.817,2.061575207e-21
252.769762,-41.772015,2507.2047,5.818413326e-21
252.647252,-41.689266,1892.3375,1.744334961e-18
252.605142,-41.641506,2249.3357,1.439763106e-20
252.938606,-41.827171,2083.404,8.544671297e-21
252.574987,-41.404007,2531.1597,6.197754090e-21
252.823988,-41.808983,5269.1013,1.717986157e-21
252.808964,-41.806942,3291.9456,3.132371216e-21
252.684978,-41.747162,5669.3965,6.109953405e-19
252.767382,-41.797062,2620.9918,5.745237922e-20
252.614036,-41.691898,2927.4283,2.601929453e-21
252.853972,-41.168407,4077.801,8.578331734e-21
252.530657,-41.494858,3710.4481,2.107929249e-21
252.888238,-41.156219,3051.514,2.338169331e-21
252.954428,-41.854729,2352.1781,7.655634184e-21
252.529755,-41.570686,3971.5989,1.646753947e-21
252.859513,-41.845554,1960.8674,2.707734881e-20
252.534145,-41.398724,2641.718,2.254402787e-20
252.517384,-41.442089,4700.5385,1.336052368e-21
252.519661,-41.420567,2712.4032,9.476834686e-21
252.505390,-41.523567,2424.7224,9.380974682e-21
252.551706,-41.664986,4533.3033,1.946284528e-21
252.533541,-41.361877,3241.7869,2.089844880e-21
252.528752,-41.363529,2884.1851,3.573968024e-21
252.537582,-41.341946,3474.6831,2.150638680e-21
252.985009,-41.116631,5300.3635,7.512473721e-21
252.943084,-41.116467,2563.0673,1.199197866e-20
252.481819,-41.505699,4124.1164,3.504032950e-21
252.684636,-41.811104,1739.5706,5.528665043e-20
252.865487,-41.120384,3887.7164,3.077710553e-21
252.625936,-41.221676,3127.7953,7.107185963e-21
252.507862,-41.351349,3286.9996,3.549705502e-21
252.479083,-41.589634,3738.6545,4.129429108e-21
252.757078,-41.859245,2270.8246,1.094835644e-20
252.554006,-41.272038,2492.9807,9.889697627e-21
252.471835,-41.598076,2736.3796,2.894022990e-21
252.542699,-41.276367,2475.8087,4.842431050e-21
252.511061,-41.312374,2496.4218,7.691660579e-21
252.554795,-41.746521,2306.537,2.806803951e-20
252.951303,-41.914600,2055.4688,2.919622148e-20
252.553309,-41.249672,2774.7686,3.598089692e-21
252.662749,-41.836929,2559.5274,6.298053060e-20
252.586034,-41.213211,2699.3648,1.272659394e-20
252.621310,-41.184315,2339.6539,5.517846047e-20
252.484728,-41.318321,3241.4826,3.034603522e-21
252.888068,-41.920795,2172.8239,1.178319618e-20
252.569062,-41.786888,1997.9861,2.540725373e-18
252.894850,-41.926277,3119.2456,6.282643464e-21
252.754071,-41.105736,2305.87,2.977926508e-19
252.416676,-41.480968,2584.3833,6.094485261e-21
252.475692,-41.304031,2771.0305,2.562855157e-21
252.959799,-41.940510,2928.3378,7.427776553e-21
252.489676,-41.275761,3113.5553,1.991796655e-21
252.818710,-41.075542,3718.8548,2.399741427e-21
252.402479,-41.491196,3794.1333,2.538537950e-21
252.442786,-41.333641,2729.4879,1.647642206e-20
252.511390,-41.766029,3978.4711,1.041812848e-19
252.733459,-41.908119,2356.8452,1.082943201e-19
252.827830,-41.937569,2212.9815,1.296106781e-20
252.590056,-41.160686,4482.565,4.093997833e-21
252.401278,-41.399200,4947.7347,2.579535530e-21
252.595412,-41.849270,3038.6398,2.195630587e-21
252.391907,-41.590279,2639.7551,6.874491362e-21
252.435847,-41.304939,2630.2127,9.154123002e-21
252.575461,-41.844299,3515.53,2.956347615e-21
252.735374,-41.926418,3943.3366,3.432362540e-21
252.543008,-41.825787,2449.224,5.869277802e-21
252.400547,-41.350430,2435.7419,2.802492442e-20
252.659622,-41.902222,1954.5765,1.958400627e-20
252.867911,-41.967499,2249.8926,7.604165132e-21
252.578956,-41.862316,3724.618,2.923059502e-21
252.797053,-41.043320,2319.464,9.471219009e-20
252.355937,-41.465790,3244.3183,2.446655993e-21
252.938864,-41.983559,2060.1537,2.139288089e-20
252.671688,-41.079185,2498.3696,1.368865176e-20
252.423228,-41.732479,3969.7407,4.496538401e-21
252.702582,-41.938839,3157.7626,2.251809079e-21
252.375407,-41.343994,4739.481,1.916162474e-21
252.747930,-41.958191,1976.0585,6.122852446e-19
252.488086,-41.817532,3065.1878,2.808993509e-21
252.531071,-41.855091,3461.697,2.165451775e-21
252.331859,-41.508068,3800.0226,1.289083783e-21
252.612445,-41.910809,2138.9346,1.076528786e-19
252.402738,-41.734364,2696.886,3.692306502e-21
252.742900,-41.032005,2248.3131,4.081614614e-20
252.403813,-41.744774,2392.2315,1.451514066e-20
252.694351,-41.044132,2796.9119,9.050642359e-21
252.449473,-41.805851,2583.8109,3.098494032e-21
252.361706,-41.310951,3088.8084,8.035276293e-20
252.628021,-41.936123,2165.4335,3.477380972e-20
252.628527,-41.062164,3525.6072,3.936962378e-21
252.498671,-41.862766,4455.1873,2.442442995e-21
252.547860,-41.900036,1829.2787,3.981719558e-19
252.511906,-41.123329,3048.4858,4.108573930e-19
252.299158,-41.447384,2850.0395,5.187928130e-21
252.335452,-41.684437,2307.2283,1.218874146e-20
252.403999,-41.211655,3878.6746,1.600194122e-21
252.313300,-41.647385,2577.9553,3.003599833e-20
252.571358,-41.070408,2709.2416,2.709037591e-20
252.425385,-41.826496,2525.6902,1.245343644e-20
252.501255,-41.108929,4557.5035,1.496484072e-21
252.419350,-41.828045,2417.7394,6.004864748e-21
252.300735,-41.656845,3045.2599,1.457684905e-20
252.553780,-41.066082,3864.3635,2.092158836e-21
252.512797,-41.088871,2482.2352,1.276192509e-20
252.585425,-41.043270,2532.8788,5.823061653e-21
252.269348,-41.407833,4001.3672,2.274827091e-21
252.366012,-41.794201,2198.7384,1.144758302e-20
252.340150,-41.238468,2289.0177,1.385972739e-19
252.343073,-41.774361,2905.8909,3.593985182e-21
252.408843,-41.852028,1865.391,5.198877258e-19
252.310223,-41.736164,3812.174,1.810697782e-21
252.320611,-41.757877,2169.6594,2.434838338e-20
252.270439,-41.333256,2365.5406,3.237555584e-20
252.348662,-41.199169,5067.5407,6.752510875e-22
252.426072,-41.885487,5469.2862,3.586713651e-20
252.539432,-41.965832,3103.0025,2.169955244e-21
252.312190,-41.232819,2607.8035,5.138810131e-21
252.322125,-41.213959,3032.9851,5.678482975e-21
252.275448,-41.723408,2303.4076,2.588029892e-20
252.501590,-41.045898,2858.9979,4.151636299e-21
252.404604,-41.113808,2089.8494,4.744769752e-19
252.244760,-41.323631,2463.2692,4.048981099e-20
252.533988,-41.983418,3385.5453,2.255859726e-21
252.277093,-41.756138,2616.2854,1.036789486e-20
252.212942,-41.389023,4550.4841,5.494747328e-21
252.327126,-41.833012,2606.1897,3.581718507e-20
252.192324,-41.478523,3399.6546,9.312523645e-21
252.195062,-41.580227,2432.3018,2.558664360e-20
252.351414,-41.132088,4835.2654,1.494196790e-20
252.219672,-41.316750,2757.7871,9.722954650e-21
252.179320,-41.542236,2505.6081,4.079715741e-20
252.384267,-41.089497,3098.6822,2.580183220e-21
252.286700,-41.818886,2391.4591,4.608540788e-21
252.404791,-41.065308,2725.8265,1.280171558e-20
252.348719,-41.109009,2563.7877,5.295587432e-21
252.275590,-41.183922,3136.0759,2.155352512e-21
252.296743,-41.850468,2642.1889,3.444749934e-21
252.155483,-41.546150,3067.0239,1.010821076e-20
252.425040,-41.032890,2191.7822,8.602034164e-20
252.257575,-41.187611,2620.5985,5.430920820e-21
252.209675,-41.750561,5933.4021,3.419106040e-21
252.139858,-41.511059,2304.5709,1.521337795e-20
252.280242,-41.859871,2554.1495,4.131733496e-20
252.290455,-41.875862,2252.1922,7.143409889e-21
252.207488,-41.773079,2741.2179,3.591469643e-21
252.268957,-41.861740,2870.4585,5.142293941e-21
252.164698,-41.705589,3226.7965,1.516741376e-21
252.136197,-41.370609,2639.5238,7.139074809e-21
252.208556,-41.799927,2404.5217,5.792148473e-21
252.121883,-41.400883,3451.9125,2.047169827e-21
252.111395,-41.560272,2620.36,2.629143700e-20
252.349896,-41.040703,2944.3872,3.233908544e-21
252.349523,-41.037327,2520.8525,7.111428303e-21
252.199753,-41.190865,4018.3583,1.468388429e-21
252.231351,-41.861416,2142.4034,3.580526979e-20
252.200156,-41.826633,2332.3668,1.323691175e-20
252.294760,-41.936317,2254.9961,8.466430175e-21
252.135159,-41.276543,2473.5397,1.114524356e-20
252.304821,-41.050865,4299.2924,1.154323215e-21
252.085210,-41.408707,2367.3699,1.484933061e-20
252.095827,-41.350533,2977.7993,2.310733205e-21
252.099867,-41.679367,4428.1414,1.260539465e-21
252.304979,-41.969177,2291.2197,7.570805420e-21
252.290798,-41.043411,4025.1584,1.694311145e-21
252.075776,-41.634266,2381.1205,8.869901372e-21
252.087812,-41.685959,4396.0093,1.769041975e-21
252.108782,-41.261692,2633.8123,3.379419937e-21
252.052929,-41.432144,2278.3046,7.148847792e-21
252.064960,-41.651070,2356.5411,1.168849360e-20
252.223106,-41.081467,2628.713,4.601469961e-21
252.232576,-41.937927,2617.5219,5.243548616e-21
252.190777,-41.899170,2013.1654,1.602521185e-20
252.081125,-41.269642,3819.9146,2.080052256e-21
252.210905,-41.930817,3236.2152,1.756321514e-21
252.035556,-41.629017,3964.5855,3.139396056e-21
252.111531,-41.189316,3886.2667,2.563179942e-21
252.035872,-41.348522,2956.2451,4.390503415e-21
252.085910,-41.218925,2288.183,6.623060966e-21
252.092595,-41.806648,3658.6782,1.857470436e-21
252.228109,-41.975403,2496.3252,1.669613407e-20
252.000450,-41.467415,2395.9855,4.911240002e-20
252.126775,-41.876106,2548.3508,8.029098280e-21
252.053633,-41.767536,2823.2556,2.497068082e-20
252.058645,-41.221397,1925.1259,1.985755782e-18
252.184620,-41.960602,2346.3856,1.004743319e-19
252.205453,-41.986469,2625.4111,6.558990736e-19
252.067377,-41.178082,2366.0375,1.062986514e-19
252.099065,-41.883869,2206.4486,1.462323794e-20
252.127616,-41.079384,5637.2404,2.605712833e-21
252.157563,-41.968189,3170.3861,3.110823479e-21
252.157831,-41.028790,3872.4118,1.005513251e-21
252.023196,-41.195976,2897.4044,3.948078463e-21
252.093282,-41.926586,2197.9598,2.684248811e-20
252.022030,-41.167690,2150.1569,1.986315993e-20
252.083875,-41.938160,1984.6353,1.319554431e-19
252.010117,-41.854439,3121.079,2.461648273e-21
252.015535,-41.126656,3400.0499,1.883892944e-21
252.010436,-41.114506,2878.126,3.049867123e-21
252.045690,-41.046795,4846.7315,2.780565962e-21
252.006986,-41.942474,4241.3759,1.703940942e-21
1 ra_deg dec_deg temperature_K amplitude_sr
2 252.995803 -41.503754 2122.6076 8.520470492e-21
3 252.993551 -41.543816 2114.6385 9.967579356e-21
4 252.918388 -41.510990 2208.128 1.446557073e-19
5 252.900742 -41.506279 2279.2582 1.973490936e-18
6 252.884436 -41.488571 2138.1302 1.410128710e-20
7 252.986092 -41.597309 2129.5884 3.136569365e-20
8 252.984711 -41.606659 2114.2229 9.627036994e-21
9 252.924427 -41.399944 3134.4748 2.266711014e-21
10 252.938142 -41.386372 5574.2493 7.758286100e-21
11 252.958086 -41.625877 1617.2207 3.144977882e-19
12 252.970634 -41.365410 2202.8994 1.648245643e-20
13 252.810573 -41.475826 3941.0552 1.524175586e-21
14 252.854571 -41.396477 2408.7178 2.646847276e-20
15 252.974001 -41.655529 2217.0089 4.616359171e-20
16 252.783806 -41.501965 2321.7548 7.149940580e-21
17 252.782782 -41.459801 2703.1671 4.949843960e-20
18 252.880635 -41.648254 2158.5784 8.959667114e-21
19 252.832964 -41.372025 2223.318 6.324985653e-20
20 252.781331 -41.583836 4289.9357 3.687496634e-21
21 252.774616 -41.415287 2503.0345 3.528731590e-20
22 252.781981 -41.604404 4481.0767 4.210660307e-21
23 252.835709 -41.344860 2471.4445 7.859657597e-20
24 252.755037 -41.586941 2520.0282 9.071531178e-21
25 252.829319 -41.663242 2092.7219 9.439286836e-21
26 252.729244 -41.560253 3870.9986 2.202578638e-21
27 252.792426 -41.349979 2347.9017 6.933492283e-21
28 252.715858 -41.556286 3197.6305 2.002903960e-21
29 252.745823 -41.618793 2093.3838 1.798894237e-20
30 252.869507 -41.293640 2565.0983 7.621517442e-21
31 252.968856 -41.268589 2990.034 2.505403621e-21
32 252.925736 -41.270466 3338.3343 2.122223459e-21
33 252.680916 -41.525299 1724.5512 7.687633387e-18
34 252.701932 -41.401005 2431.4439 5.080775898e-20
35 252.829047 -41.712399 2045.1192 5.651479174e-20
36 252.678461 -41.426064 3910.3355 1.929630882e-21
37 252.944337 -41.247566 2738.4573 2.298467128e-20
38 252.821892 -41.277275 2590.152 3.408118915e-21
39 252.662337 -41.573349 2041.2244 1.499212709e-20
40 252.651190 -41.553684 3393.4186 2.240876206e-21
41 252.717832 -41.332279 2490.4613 6.309850262e-20
42 252.859629 -41.752144 1851.5385 1.237866868e-19
43 252.710481 -41.329098 3931.3812 2.005228197e-21
44 252.712726 -41.321545 3857.1544 5.275918149e-21
45 252.907975 -41.225502 2783.4349 7.275926125e-21
46 252.617463 -41.497688 2452.5593 1.840842057e-20
47 252.697290 -41.680248 2279.1812 1.329110424e-20
48 252.701091 -41.312347 2287.752 8.783238970e-21
49 252.641947 -41.624428 2447.7312 1.460562478e-20
50 252.611005 -41.565239 2414.9482 7.249220492e-21
51 252.921863 -41.204163 2880.0791 4.126485206e-21
52 252.639833 -41.358757 3812.7671 1.743956802e-21
53 252.594286 -41.548817 3760.853 2.953969609e-21
54 252.630135 -41.641548 2306.0382 3.363985170e-19
55 252.825565 -41.215195 2246.365 7.392177336e-20
56 252.620751 -41.361652 2520.4538 1.047956906e-20
57 252.802403 -41.217541 4186.817 2.061575207e-21
58 252.769762 -41.772015 2507.2047 5.818413326e-21
59 252.647252 -41.689266 1892.3375 1.744334961e-18
60 252.605142 -41.641506 2249.3357 1.439763106e-20
61 252.938606 -41.827171 2083.404 8.544671297e-21
62 252.574987 -41.404007 2531.1597 6.197754090e-21
63 252.823988 -41.808983 5269.1013 1.717986157e-21
64 252.808964 -41.806942 3291.9456 3.132371216e-21
65 252.684978 -41.747162 5669.3965 6.109953405e-19
66 252.767382 -41.797062 2620.9918 5.745237922e-20
67 252.614036 -41.691898 2927.4283 2.601929453e-21
68 252.853972 -41.168407 4077.801 8.578331734e-21
69 252.530657 -41.494858 3710.4481 2.107929249e-21
70 252.888238 -41.156219 3051.514 2.338169331e-21
71 252.954428 -41.854729 2352.1781 7.655634184e-21
72 252.529755 -41.570686 3971.5989 1.646753947e-21
73 252.859513 -41.845554 1960.8674 2.707734881e-20
74 252.534145 -41.398724 2641.718 2.254402787e-20
75 252.517384 -41.442089 4700.5385 1.336052368e-21
76 252.519661 -41.420567 2712.4032 9.476834686e-21
77 252.505390 -41.523567 2424.7224 9.380974682e-21
78 252.551706 -41.664986 4533.3033 1.946284528e-21
79 252.533541 -41.361877 3241.7869 2.089844880e-21
80 252.528752 -41.363529 2884.1851 3.573968024e-21
81 252.537582 -41.341946 3474.6831 2.150638680e-21
82 252.985009 -41.116631 5300.3635 7.512473721e-21
83 252.943084 -41.116467 2563.0673 1.199197866e-20
84 252.481819 -41.505699 4124.1164 3.504032950e-21
85 252.684636 -41.811104 1739.5706 5.528665043e-20
86 252.865487 -41.120384 3887.7164 3.077710553e-21
87 252.625936 -41.221676 3127.7953 7.107185963e-21
88 252.507862 -41.351349 3286.9996 3.549705502e-21
89 252.479083 -41.589634 3738.6545 4.129429108e-21
90 252.757078 -41.859245 2270.8246 1.094835644e-20
91 252.554006 -41.272038 2492.9807 9.889697627e-21
92 252.471835 -41.598076 2736.3796 2.894022990e-21
93 252.542699 -41.276367 2475.8087 4.842431050e-21
94 252.511061 -41.312374 2496.4218 7.691660579e-21
95 252.554795 -41.746521 2306.537 2.806803951e-20
96 252.951303 -41.914600 2055.4688 2.919622148e-20
97 252.553309 -41.249672 2774.7686 3.598089692e-21
98 252.662749 -41.836929 2559.5274 6.298053060e-20
99 252.586034 -41.213211 2699.3648 1.272659394e-20
100 252.621310 -41.184315 2339.6539 5.517846047e-20
101 252.484728 -41.318321 3241.4826 3.034603522e-21
102 252.888068 -41.920795 2172.8239 1.178319618e-20
103 252.569062 -41.786888 1997.9861 2.540725373e-18
104 252.894850 -41.926277 3119.2456 6.282643464e-21
105 252.754071 -41.105736 2305.87 2.977926508e-19
106 252.416676 -41.480968 2584.3833 6.094485261e-21
107 252.475692 -41.304031 2771.0305 2.562855157e-21
108 252.959799 -41.940510 2928.3378 7.427776553e-21
109 252.489676 -41.275761 3113.5553 1.991796655e-21
110 252.818710 -41.075542 3718.8548 2.399741427e-21
111 252.402479 -41.491196 3794.1333 2.538537950e-21
112 252.442786 -41.333641 2729.4879 1.647642206e-20
113 252.511390 -41.766029 3978.4711 1.041812848e-19
114 252.733459 -41.908119 2356.8452 1.082943201e-19
115 252.827830 -41.937569 2212.9815 1.296106781e-20
116 252.590056 -41.160686 4482.565 4.093997833e-21
117 252.401278 -41.399200 4947.7347 2.579535530e-21
118 252.595412 -41.849270 3038.6398 2.195630587e-21
119 252.391907 -41.590279 2639.7551 6.874491362e-21
120 252.435847 -41.304939 2630.2127 9.154123002e-21
121 252.575461 -41.844299 3515.53 2.956347615e-21
122 252.735374 -41.926418 3943.3366 3.432362540e-21
123 252.543008 -41.825787 2449.224 5.869277802e-21
124 252.400547 -41.350430 2435.7419 2.802492442e-20
125 252.659622 -41.902222 1954.5765 1.958400627e-20
126 252.867911 -41.967499 2249.8926 7.604165132e-21
127 252.578956 -41.862316 3724.618 2.923059502e-21
128 252.797053 -41.043320 2319.464 9.471219009e-20
129 252.355937 -41.465790 3244.3183 2.446655993e-21
130 252.938864 -41.983559 2060.1537 2.139288089e-20
131 252.671688 -41.079185 2498.3696 1.368865176e-20
132 252.423228 -41.732479 3969.7407 4.496538401e-21
133 252.702582 -41.938839 3157.7626 2.251809079e-21
134 252.375407 -41.343994 4739.481 1.916162474e-21
135 252.747930 -41.958191 1976.0585 6.122852446e-19
136 252.488086 -41.817532 3065.1878 2.808993509e-21
137 252.531071 -41.855091 3461.697 2.165451775e-21
138 252.331859 -41.508068 3800.0226 1.289083783e-21
139 252.612445 -41.910809 2138.9346 1.076528786e-19
140 252.402738 -41.734364 2696.886 3.692306502e-21
141 252.742900 -41.032005 2248.3131 4.081614614e-20
142 252.403813 -41.744774 2392.2315 1.451514066e-20
143 252.694351 -41.044132 2796.9119 9.050642359e-21
144 252.449473 -41.805851 2583.8109 3.098494032e-21
145 252.361706 -41.310951 3088.8084 8.035276293e-20
146 252.628021 -41.936123 2165.4335 3.477380972e-20
147 252.628527 -41.062164 3525.6072 3.936962378e-21
148 252.498671 -41.862766 4455.1873 2.442442995e-21
149 252.547860 -41.900036 1829.2787 3.981719558e-19
150 252.511906 -41.123329 3048.4858 4.108573930e-19
151 252.299158 -41.447384 2850.0395 5.187928130e-21
152 252.335452 -41.684437 2307.2283 1.218874146e-20
153 252.403999 -41.211655 3878.6746 1.600194122e-21
154 252.313300 -41.647385 2577.9553 3.003599833e-20
155 252.571358 -41.070408 2709.2416 2.709037591e-20
156 252.425385 -41.826496 2525.6902 1.245343644e-20
157 252.501255 -41.108929 4557.5035 1.496484072e-21
158 252.419350 -41.828045 2417.7394 6.004864748e-21
159 252.300735 -41.656845 3045.2599 1.457684905e-20
160 252.553780 -41.066082 3864.3635 2.092158836e-21
161 252.512797 -41.088871 2482.2352 1.276192509e-20
162 252.585425 -41.043270 2532.8788 5.823061653e-21
163 252.269348 -41.407833 4001.3672 2.274827091e-21
164 252.366012 -41.794201 2198.7384 1.144758302e-20
165 252.340150 -41.238468 2289.0177 1.385972739e-19
166 252.343073 -41.774361 2905.8909 3.593985182e-21
167 252.408843 -41.852028 1865.391 5.198877258e-19
168 252.310223 -41.736164 3812.174 1.810697782e-21
169 252.320611 -41.757877 2169.6594 2.434838338e-20
170 252.270439 -41.333256 2365.5406 3.237555584e-20
171 252.348662 -41.199169 5067.5407 6.752510875e-22
172 252.426072 -41.885487 5469.2862 3.586713651e-20
173 252.539432 -41.965832 3103.0025 2.169955244e-21
174 252.312190 -41.232819 2607.8035 5.138810131e-21
175 252.322125 -41.213959 3032.9851 5.678482975e-21
176 252.275448 -41.723408 2303.4076 2.588029892e-20
177 252.501590 -41.045898 2858.9979 4.151636299e-21
178 252.404604 -41.113808 2089.8494 4.744769752e-19
179 252.244760 -41.323631 2463.2692 4.048981099e-20
180 252.533988 -41.983418 3385.5453 2.255859726e-21
181 252.277093 -41.756138 2616.2854 1.036789486e-20
182 252.212942 -41.389023 4550.4841 5.494747328e-21
183 252.327126 -41.833012 2606.1897 3.581718507e-20
184 252.192324 -41.478523 3399.6546 9.312523645e-21
185 252.195062 -41.580227 2432.3018 2.558664360e-20
186 252.351414 -41.132088 4835.2654 1.494196790e-20
187 252.219672 -41.316750 2757.7871 9.722954650e-21
188 252.179320 -41.542236 2505.6081 4.079715741e-20
189 252.384267 -41.089497 3098.6822 2.580183220e-21
190 252.286700 -41.818886 2391.4591 4.608540788e-21
191 252.404791 -41.065308 2725.8265 1.280171558e-20
192 252.348719 -41.109009 2563.7877 5.295587432e-21
193 252.275590 -41.183922 3136.0759 2.155352512e-21
194 252.296743 -41.850468 2642.1889 3.444749934e-21
195 252.155483 -41.546150 3067.0239 1.010821076e-20
196 252.425040 -41.032890 2191.7822 8.602034164e-20
197 252.257575 -41.187611 2620.5985 5.430920820e-21
198 252.209675 -41.750561 5933.4021 3.419106040e-21
199 252.139858 -41.511059 2304.5709 1.521337795e-20
200 252.280242 -41.859871 2554.1495 4.131733496e-20
201 252.290455 -41.875862 2252.1922 7.143409889e-21
202 252.207488 -41.773079 2741.2179 3.591469643e-21
203 252.268957 -41.861740 2870.4585 5.142293941e-21
204 252.164698 -41.705589 3226.7965 1.516741376e-21
205 252.136197 -41.370609 2639.5238 7.139074809e-21
206 252.208556 -41.799927 2404.5217 5.792148473e-21
207 252.121883 -41.400883 3451.9125 2.047169827e-21
208 252.111395 -41.560272 2620.36 2.629143700e-20
209 252.349896 -41.040703 2944.3872 3.233908544e-21
210 252.349523 -41.037327 2520.8525 7.111428303e-21
211 252.199753 -41.190865 4018.3583 1.468388429e-21
212 252.231351 -41.861416 2142.4034 3.580526979e-20
213 252.200156 -41.826633 2332.3668 1.323691175e-20
214 252.294760 -41.936317 2254.9961 8.466430175e-21
215 252.135159 -41.276543 2473.5397 1.114524356e-20
216 252.304821 -41.050865 4299.2924 1.154323215e-21
217 252.085210 -41.408707 2367.3699 1.484933061e-20
218 252.095827 -41.350533 2977.7993 2.310733205e-21
219 252.099867 -41.679367 4428.1414 1.260539465e-21
220 252.304979 -41.969177 2291.2197 7.570805420e-21
221 252.290798 -41.043411 4025.1584 1.694311145e-21
222 252.075776 -41.634266 2381.1205 8.869901372e-21
223 252.087812 -41.685959 4396.0093 1.769041975e-21
224 252.108782 -41.261692 2633.8123 3.379419937e-21
225 252.052929 -41.432144 2278.3046 7.148847792e-21
226 252.064960 -41.651070 2356.5411 1.168849360e-20
227 252.223106 -41.081467 2628.713 4.601469961e-21
228 252.232576 -41.937927 2617.5219 5.243548616e-21
229 252.190777 -41.899170 2013.1654 1.602521185e-20
230 252.081125 -41.269642 3819.9146 2.080052256e-21
231 252.210905 -41.930817 3236.2152 1.756321514e-21
232 252.035556 -41.629017 3964.5855 3.139396056e-21
233 252.111531 -41.189316 3886.2667 2.563179942e-21
234 252.035872 -41.348522 2956.2451 4.390503415e-21
235 252.085910 -41.218925 2288.183 6.623060966e-21
236 252.092595 -41.806648 3658.6782 1.857470436e-21
237 252.228109 -41.975403 2496.3252 1.669613407e-20
238 252.000450 -41.467415 2395.9855 4.911240002e-20
239 252.126775 -41.876106 2548.3508 8.029098280e-21
240 252.053633 -41.767536 2823.2556 2.497068082e-20
241 252.058645 -41.221397 1925.1259 1.985755782e-18
242 252.184620 -41.960602 2346.3856 1.004743319e-19
243 252.205453 -41.986469 2625.4111 6.558990736e-19
244 252.067377 -41.178082 2366.0375 1.062986514e-19
245 252.099065 -41.883869 2206.4486 1.462323794e-20
246 252.127616 -41.079384 5637.2404 2.605712833e-21
247 252.157563 -41.968189 3170.3861 3.110823479e-21
248 252.157831 -41.028790 3872.4118 1.005513251e-21
249 252.023196 -41.195976 2897.4044 3.948078463e-21
250 252.093282 -41.926586 2197.9598 2.684248811e-20
251 252.022030 -41.167690 2150.1569 1.986315993e-20
252 252.083875 -41.938160 1984.6353 1.319554431e-19
253 252.010117 -41.854439 3121.079 2.461648273e-21
254 252.015535 -41.126656 3400.0499 1.883892944e-21
255 252.010436 -41.114506 2878.126 3.049867123e-21
256 252.045690 -41.046795 4846.7315 2.780565962e-21
257 252.006986 -41.942474 4241.3759 1.703940942e-21
benchmarks/hip_psf_2026-09-07/catalog_8192.csv
+8193
View File
File diff suppressed because it is too large. Load diff
benchmarks/hip_psf_2026-09-07/clean_baseline_multi.log
+10
View File
@@ -0,0 +1,10 @@
PSF cache ready: 64x64 phases, radius 47 px, relative tail 1e-08, tail abs 1e-06, boundary 1e-07, build 0.689 s
Frame 0: finding and prefetching catalog tiles...
Catalog prefetch: 64440 requested, 3 newly loaded (21087 stars), 64437 unavailable in 0.033 s; 4 loader workers
Frame 0: catalog prefetch finished (64440 candidate tiles).
Frame 0: splatting 130882 lens triangles...
Frame 0: catalog splatting finished in 1.9 s; writing image...
Rendered 156837 images from 21087 catalog stars to /tmp/gr-hip-clean-baseline-multi.png (ok; imported lens map)
PSF splats: cached 156837, cached wing-clipped 0, direct fallbacks 0, discarded below min-Y 0
HIP PSF: 156837 events in 10 batches (10 timed); upload 0.000304 s, kernel 1.496493 s, download 0.026518 s
OMP_NUM_THREADS=16 timeout --kill-after=5s 60s --all-sky-catalog 4 1e-8 12.48s user 0.25s system 421% cpu 3.016 total
benchmarks/hip_psf_2026-09-07/cpu_multi.log
+209
View File
@@ -0,0 +1,209 @@
PSF cache ready: 64x64 phases, radius 47 px, relative tail 1e-08, tail abs 1e-06, boundary 1e-07, build 0.662 s
Frame 0: finding and prefetching catalog tiles...
Catalog prefetch: 64440 requested, 3 newly loaded (21087 stars), 64437 unavailable in 0.035 s; 4 loader workers
Frame 0: catalog prefetch finished (64440 candidate tiles).
Frame 0: splatting 130882 lens triangles...
Frame 0: splat worker 2/16 started.
Frame 0: splat worker 4/16 started.
Frame 0: splat worker 3/16 started.
Frame 0: splat worker 3/16 reached 8 local triangles.
Frame 0: splat worker 1/16 started.
Frame 0: splat worker 4/16 reached 8 local triangles.
Frame 0: splat worker 4/16 reached 16 local triangles.
Frame 0: splat worker 2/16 reached 8 local triangles.
Frame 0: splat worker 1/16 reached 8 local triangles.
Frame 0: splat worker 7/16 started.
Frame 0: splat worker 4/16 reached 32 local triangles.
Frame 0: splat worker 6/16 started.
Frame 0: splat worker 2/16 reached 16 local triangles.
Frame 0: splat worker 1/16 reached 16 local triangles.
Frame 0: splat worker 3/16 reached 16 local triangles.
Frame 0: splat worker 7/16 reached 8 local triangles.
Frame 0: splat worker 4/16 reached 64 local triangles.
Frame 0: splat worker 6/16 reached 8 local triangles.
Frame 0: splat worker 2/16 reached 32 local triangles.
Frame 0: splat worker 1/16 reached 32 local triangles.
Frame 0: splat worker 3/16 reached 32 local triangles.
Frame 0: splat worker 7/16 reached 16 local triangles.
Frame 0: splat worker 3/16 reached 64 local triangles.
Frame 0: splat worker 4/16 reached 128 local triangles.
Frame 0: splat worker 2/16 reached 64 local triangles.
Frame 0: splat worker 1/16 reached 64 local triangles.
Frame 0: splat worker 3/16 reached 128 local triangles.
Frame 0: splat worker 6/16 reached 16 local triangles.
Frame 0: splat worker 7/16 reached 32 local triangles.
Frame 0: splat worker 6/16 reached 32 local triangles.
Frame 0: splat worker 14/16 started.
Frame 0: splat worker 15/16 started.
Frame 0: splat worker 13/16 started.
Frame 0: splat worker 12/16 started.
Frame 0: splat worker 11/16 started.
Frame 0: splat worker 16/16 started.
Frame 0: splat worker 2/16 reached 128 local triangles.
Frame 0: splat worker 10/16 started.
Frame 0: splat worker 9/16 started.
Frame 0: splat worker 8/16 started.
Frame 0: splat worker 15/16 reached 8 local triangles.
Frame 0: splat worker 1/16 reached 128 local triangles.
Frame 0: splat worker 15/16 reached 16 local triangles.
Frame 0: splat worker 9/16 reached 8 local triangles.
Frame 0: splat worker 12/16 reached 8 local triangles.
Frame 0: splat worker 10/16 reached 8 local triangles.
Frame 0: splat worker 10/16 reached 16 local triangles.
Frame 0: splat worker 1/16 reached 256 local triangles.
Frame 0: splat worker 10/16 reached 32 local triangles.
Frame 0: splat worker 12/16 reached 16 local triangles.
Frame 0: splat worker 5/16 started.
Frame 0: splat worker 12/16 reached 32 local triangles.
Frame 0: splat worker 5/16 reached 8 local triangles.
Frame 0: splat worker 9/16 reached 16 local triangles.
Frame 0: splat worker 5/16 reached 16 local triangles.
Frame 0: splat worker 10/16 reached 64 local triangles.
Frame 0: splat worker 15/16 reached 32 local triangles.
Frame 0: splat worker 13/16 reached 8 local triangles.
Frame 0: splat worker 13/16 reached 16 local triangles.
Frame 0: splat worker 13/16 reached 32 local triangles.
Frame 0: splat worker 13/16 reached 64 local triangles.
Frame 0: splat worker 7/16 reached 64 local triangles.
Frame 0: splat worker 6/16 reached 64 local triangles.
Frame 0: splat worker 13/16 reached 128 local triangles.
Frame 0: splat worker 4/16 reached 256 local triangles.
Frame 0: splat worker 7/16 reached 128 local triangles.
Frame 0: splat worker 14/16 reached 8 local triangles.
Frame 0: splat worker 11/16 reached 8 local triangles.
Frame 0: splat worker 12/16 reached 64 local triangles.
Frame 0: splat worker 2/16 reached 256 local triangles.
Frame 0: splat worker 9/16 reached 32 local triangles.
Frame 0: splat worker 8/16 reached 8 local triangles.
Frame 0: splat worker 14/16 reached 16 local triangles.
Frame 0: splat worker 11/16 reached 16 local triangles.
Frame 0: splat worker 5/16 reached 32 local triangles.
Frame 0: splat worker 8/16 reached 16 local triangles.
Frame 0: splat worker 6/16 reached 128 local triangles.
Frame 0: splat worker 14/16 reached 32 local triangles.
Frame 0: splat worker 16/16 reached 8 local triangles.
Frame 0: splat worker 12/16 reached 128 local triangles.
Frame 0: splat worker 13/16 reached 256 local triangles.
Frame 0: splat worker 8/16 reached 32 local triangles.
Frame 0: splat worker 11/16 reached 32 local triangles.
Frame 0: splat worker 15/16 reached 64 local triangles.
Frame 0: splat worker 9/16 reached 64 local triangles.
Frame 0: splat worker 16/16 reached 16 local triangles.
Frame 0: splat worker 5/16 reached 64 local triangles.
Frame 0: splat worker 8/16 reached 64 local triangles.
Frame 0: splat worker 8/16 reached 128 local triangles.
Frame 0: splat worker 9/16 reached 128 local triangles.
Frame 0: splat worker 8/16 reached 256 local triangles.
Frame 0: splat worker 16/16 reached 32 local triangles.
Frame 0: splat worker 14/16 reached 64 local triangles.
Frame 0: splat worker 11/16 reached 64 local triangles.
Frame 0: splat worker 9/16 reached 256 local triangles.
Frame 0: splat worker 1/16 reached 512 local triangles.
Frame 0: splat worker 7/16 reached 256 local triangles.
Frame 0: splat worker 3/16 reached 256 local triangles.
Frame 0: splat worker 14/16 reached 128 local triangles.
Frame 0: splat worker 4/16 reached 512 local triangles.
Frame 0: splat worker 5/16 reached 128 local triangles.
Frame 0: splat worker 14/16 reached 256 local triangles.
Frame 0: splat worker 15/16 reached 128 local triangles.
Frame 0: splat worker 2/16 reached 512 local triangles.
Frame 0: splat worker 6/16 reached 256 local triangles.
Frame 0: splat worker 5/16 reached 256 local triangles.
Frame 0: splat worker 16/16 reached 64 local triangles.
Frame 0: splat worker 11/16 reached 128 local triangles.
Frame 0: splat worker 3/16 reached 512 local triangles.
Frame 0: splat worker 12/16 reached 256 local triangles.
Frame 0: splat worker 13/16 reached 512 local triangles.
Frame 0: splat worker 10/16 reached 128 local triangles.
Frame 0: splat worker 8/16 reached 512 local triangles.
Frame 0: splat worker 10/16 reached 256 local triangles.
Frame 0: splat worker 6/16 reached 512 local triangles.
Frame 0: splat worker 14/16 reached 512 local triangles.
Frame 0: splat worker 9/16 reached 512 local triangles.
Frame 0: splat worker 1/16 reached 1024 local triangles.
Frame 0: splat worker 4/16 reached 1024 local triangles.
Frame 0: splat worker 13/16 reached 1024 local triangles.
Frame 0: splat worker 11/16 reached 256 local triangles.
Frame 0: splat worker 5/16 reached 512 local triangles.
Frame 0: splat worker 15/16 reached 256 local triangles.
Frame 0: splat worker 16/16 reached 128 local triangles.
Frame 0: splat worker 7/16 reached 512 local triangles.
Frame 0: splat worker 12/16 reached 512 local triangles.
Frame 0: splat worker 15/16 reached 512 local triangles.
Frame 0: splat worker 16/16 reached 256 local triangles.
Frame 0: splat worker 11/16 reached 512 local triangles.
Frame 0: splat worker 6/16 reached 1024 local triangles.
Frame 0: splat worker 3/16 reached 1024 local triangles.
Frame 0: splat worker 2/16 reached 1024 local triangles.
Frame 0: splat worker 7/16 reached 1024 local triangles.
Frame 0: splat worker 8/16 reached 1024 local triangles.
Frame 0: splat worker 10/16 reached 512 local triangles.
Frame 0: splat worker 14/16 reached 1024 local triangles.
Frame 0: splat worker 13/16 reached 2048 local triangles.
Frame 0: splat worker 12/16 reached 1024 local triangles.
Frame 0: splat worker 5/16 reached 1024 local triangles.
Frame 0: splat worker 15/16 reached 1024 local triangles.
Frame 0: splat worker 9/16 reached 1024 local triangles.
Frame 0: splat worker 4/16 reached 2048 local triangles.
Frame 0: splat worker 11/16 reached 1024 local triangles.
Frame 0: splat worker 3/16 reached 2048 local triangles.
Frame 0: splat worker 16/16 reached 512 local triangles.
Frame 0: splat worker 14/16 reached 2048 local triangles.
Frame 0: splat worker 1/16 reached 2048 local triangles.
Frame 0: splat worker 10/16 reached 1024 local triangles.
Frame 0: splat worker 2/16 reached 2048 local triangles.
Frame 0: splat worker 8/16 reached 2048 local triangles.
Frame 0: splat worker 15/16 reached 2048 local triangles.
Frame 0: splat worker 11/16 reached 2048 local triangles.
Frame 0: splat worker 5/16 reached 2048 local triangles.
Frame 0: splat worker 9/16 reached 2048 local triangles.
Frame 0: splat worker 6/16 reached 2048 local triangles.
Frame 0: splat worker 13/16 reached 4096 local triangles.
Frame 0: splat worker 12/16 reached 2048 local triangles.
Frame 0: splat worker 7/16 reached 2048 local triangles.
Frame 0: splat worker 3/16 reached 4096 local triangles.
Frame 0: splat worker 4/16 reached 4096 local triangles.
Frame 0: splat worker 8/16 reached 4096 local triangles.
Frame 0: splat worker 14/16 reached 4096 local triangles.
Frame 0: splat worker 6/16 reached 4096 local triangles.
Frame 0: splat worker 15/16 reached 4096 local triangles.
Frame 0: splat worker 16/16 reached 1024 local triangles.
Frame 0: splat worker 5/16 reached 4096 local triangles.
Frame 0: splat worker 1/16 reached 4096 local triangles.
Frame 0: splat worker 10/16 reached 2048 local triangles.
Frame 0: splat worker 13/16 reached 8192 local triangles.
Frame 0: splat worker 7/16 reached 4096 local triangles.
Frame 0: splat worker 16/16 reached 2048 local triangles.
Frame 0: splat worker 9/16 reached 4096 local triangles.
Frame 0: splat worker 12/16 reached 4096 local triangles.
Frame 0: splat worker 11/16 reached 4096 local triangles.
Frame 0: splat worker 2/16 reached 4096 local triangles.
Frame 0: splat worker 3/16 reached 8192 local triangles.
Frame 0: splat worker 14/16 reached 8192 local triangles.
Frame 0: splat worker 15/16 reached 8192 local triangles.
Frame 0: splat worker 6/16 reached 8192 local triangles.
Frame 0: splat worker 4/16 reached 8192 local triangles.
Frame 0: splat worker 1/16 reached 8192 local triangles.
Frame 0: splat worker 16/16 reached 4096 local triangles.
Frame 0: splat worker 10/16 reached 4096 local triangles.
Frame 0: splat worker 8/16 reached 8192 local triangles.
Frame 0: splat worker 3/16 finished after 10458 local triangles.
Frame 0: splat worker 16/16 finished after 4920 local triangles.
Frame 0: splat worker 11/16 finished after 6992 local triangles.
Frame 0: splat worker 10/16 finished after 4986 local triangles.
Frame 0: splat worker 15/16 finished after 9498 local triangles.
Frame 0: splat worker 6/16 finished after 8930 local triangles.
Frame 0: splat worker 4/16 finished after 9108 local triangles.
Frame 0: splat worker 12/16 finished after 7689 local triangles.
Frame 0: splat worker 5/16 finished after 7943 local triangles.
Frame 0: splat worker 14/16 finished after 9623 local triangles.
Frame 0: splat worker 8/16 finished after 8856 local triangles.
Frame 0: splat worker 1/16 finished after 9031 local triangles.
Frame 0: splat worker 2/16 finished after 7194 local triangles.
Frame 0: splat worker 9/16 finished after 7002 local triangles.
Frame 0: splat worker 13/16 finished after 11036 local triangles.
Frame 0: splat worker 7/16 finished after 7616 local triangles.
Frame 0: catalog splatting finished in 0.6 s; writing image...
Rendered 156837 images from 21087 catalog stars to /tmp/gr-cpu-multi.png (ok; imported lens map)
PSF splats: cached 156837, cached wing-clipped 0, direct fallbacks 0, discarded below min-Y 0
OMP_NUM_THREADS=16 timeout --kill-after=5s 60s --all-sky-catalog 4 1e-8 17.59s user 1.25s system 1125% cpu 1.674 total
benchmarks/hip_psf_2026-09-07/cpu_regression.log
+41
View File
@@ -0,0 +1,41 @@
cc -DENABLE_PNG -DFRAME_PSF_EVENT_SINK=1 -std=c11 -march=native -pipe -Wall -Wextra -Wpedantic -O2 -DNDEBUG -fopenmp -Isrc tests/test_observer.c src/catalog.c src/frame.c src/geodesic.c src/lens_map.c src/movie.c src/observer.c src/observer_track.c src/optics.c src/ray.c src/spacetime_common.c src/spacetime_minkowski.c -lm -lpng -o build/Release/test_observer_minkowski
cc -DENABLE_PNG -DFRAME_PSF_EVENT_SINK=1 -std=c11 -march=native -pipe -Wall -Wextra -Wpedantic -O2 -DNDEBUG -fopenmp -DSPACETIME_SCHWARZSCHILD -Isrc tests/test_observer.c src/catalog.c src/frame.c src/geodesic.c src/lens_map.c src/movie.c src/observer.c src/observer_track.c src/optics.c src/ray.c src/spacetime_common.c src/spacetime_schwarzschild.c -lm -lpng -o build/Release/test_observer_schwarzschild
cc -DENABLE_PNG -DFRAME_PSF_EVENT_SINK=1 -std=c11 -march=native -pipe -Wall -Wextra -Wpedantic -O2 -DNDEBUG -fopenmp -Isrc tests/test_geodesic.c src/catalog.c src/frame.c src/geodesic.c src/lens_map.c src/movie.c src/observer.c src/observer_track.c src/optics.c src/ray.c src/spacetime_common.c src/spacetime_minkowski.c -lm -lpng -o build/Release/test_geodesic
cc -DENABLE_PNG -DFRAME_PSF_EVENT_SINK=1 -std=c11 -march=native -pipe -Wall -Wextra -Wpedantic -O2 -DNDEBUG -fopenmp -Isrc tests/test_frame.c src/catalog.c src/frame.c src/geodesic.c src/lens_map.c src/movie.c src/observer.c src/observer_track.c src/optics.c src/ray.c src/spacetime_common.c src/spacetime_minkowski.c -lm -lpng -o build/Release/test_frame
cc -DENABLE_PNG -DFRAME_PSF_EVENT_SINK=1 -std=c11 -march=native -pipe -Wall -Wextra -Wpedantic -O2 -DNDEBUG -fopenmp -Isrc tests/test_schwarzschild.c src/catalog.c src/frame.c src/geodesic.c src/lens_map.c src/movie.c src/observer.c src/observer_track.c src/optics.c src/ray.c src/spacetime_common.c src/spacetime_schwarzschild.c -lm -lpng -o build/Release/test_schwarzschild
cc -DENABLE_PNG -DFRAME_PSF_EVENT_SINK=1 -std=c11 -march=native -pipe -Wall -Wextra -Wpedantic -O2 -DNDEBUG -fopenmp -Isrc tests/test_observer_track.c src/catalog.c src/frame.c src/geodesic.c src/lens_map.c src/movie.c src/observer.c src/observer_track.c src/optics.c src/ray.c src/spacetime_common.c -lm -lpng -o build/Release/test_observer_track
cc -DENABLE_PNG -DFRAME_PSF_EVENT_SINK=1 -std=c11 -march=native -pipe -Wall -Wextra -Wpedantic -O2 -DNDEBUG -fopenmp -Isrc tests/test_catalog_prefetch.c src/catalog.c src/frame.c src/geodesic.c src/lens_map.c src/movie.c src/observer.c src/observer_track.c src/optics.c src/ray.c src/spacetime_common.c src/spacetime_minkowski.c -lm -lpng -o build/Release/test_catalog_prefetch
make SPACETIME=minkowski ENABLE_HDR=1 backend
make[1]: Entering directory '/home/wyj/Code/C/GR_4d_raytracing'
cc -O2 -DNDEBUG -fopenmp build/Release/obj/minkowski/hdr_sink1_cpu/src/catalog.o build/Release/obj/minkowski/hdr_sink1_cpu/src/frame.o build/Release/obj/minkowski/hdr_sink1_cpu/src/geodesic.o build/Release/obj/minkowski/hdr_sink1_cpu/src/lens_map.o build/Release/obj/minkowski/hdr_sink1_cpu/src/movie.o build/Release/obj/minkowski/hdr_sink1_cpu/src/observer.o build/Release/obj/minkowski/hdr_sink1_cpu/src/observer_track.o build/Release/obj/minkowski/hdr_sink1_cpu/src/optics.o build/Release/obj/minkowski/hdr_sink1_cpu/src/ray.o build/Release/obj/minkowski/hdr_sink1_cpu/src/spacetime_common.o build/Release/obj/minkowski/hdr_sink1_cpu/src/spacetime_minkowski.o build/Release/obj/minkowski/hdr_sink1_cpu/src/main.o -lm -lpng -lcfitsio -o build/Release/minkowski_sky
make[1]: Leaving directory '/home/wyj/Code/C/GR_4d_raytracing'
mkdir -p /tmp/gr_psf_event_sink_reference
OMP_NUM_THREADS=16 build/Release/minkowski_sky --catalog assets/sky_grid_5deg.csv --output /tmp/gr_psf_event_sink_reference/minkowski_ra1_dec1_640x360.png --hdr-output --width 640 --height 360 --fov-deg 30 --look-ra-deg 1 --look-dec-deg 1 --exposure 0.1 --observer-radius 30 --observer-velocity 0 0 0 --psf-fwhm-pixels 2.7 --psf-moffat-beta 4.5 --max-magnification 1e300 --max-cache-psf-flux 1 --psf-relative-tail 1e-8 --psf-min-y 0 --coarse-cell-pixels 16 --refine-max-level 0 --refine-angle-abs-deg 0.001 --refine-angle-rel 0.1 --refine-jacobian-min 1e-3 --refine-min-edge-pixels 0.5 --refine-min-area-pixels2 0.25 --catalog-load-workers 4
PSF cache ready: 64x64 phases, radius 47 px, relative tail 1e-08, tail abs 1e-06, boundary 1e-07, build 0.677 s
Rendered 35 images from 5654 catalog stars to /tmp/gr_psf_event_sink_reference/minkowski_ra1_dec1_640x360.png (ok)
PSF splats: cached 35, cached wing-clipped 0, direct fallbacks 0, discarded below min-Y 0
python3 scripts/fits_floatdiff.py tests/data/psf_event_sink_reference/minkowski_ra1_dec1_640x360_HDR.fits /tmp/gr_psf_event_sink_reference/minkowski_ra1_dec1_640x360_HDR.fits
floatdiff: shape=(640, 360, 3) samples=691200 mismatches=0 max_abs=0 max_rel=0 worst_sample=0
make SPACETIME=schwarzschild ENABLE_HDR=1 backend
make[1]: Entering directory '/home/wyj/Code/C/GR_4d_raytracing'
cc -O2 -DNDEBUG -fopenmp build/Release/obj/schwarzschild/hdr_sink1_cpu/src/catalog.o build/Release/obj/schwarzschild/hdr_sink1_cpu/src/frame.o build/Release/obj/schwarzschild/hdr_sink1_cpu/src/geodesic.o build/Release/obj/schwarzschild/hdr_sink1_cpu/src/lens_map.o build/Release/obj/schwarzschild/hdr_sink1_cpu/src/movie.o build/Release/obj/schwarzschild/hdr_sink1_cpu/src/observer.o build/Release/obj/schwarzschild/hdr_sink1_cpu/src/observer_track.o build/Release/obj/schwarzschild/hdr_sink1_cpu/src/optics.o build/Release/obj/schwarzschild/hdr_sink1_cpu/src/ray.o build/Release/obj/schwarzschild/hdr_sink1_cpu/src/spacetime_common.o build/Release/obj/schwarzschild/hdr_sink1_cpu/src/spacetime_schwarzschild.o build/Release/obj/schwarzschild/hdr_sink1_cpu/src/main.o -lm -lpng -lcfitsio -o build/Release/schwarzschild_sky
make[1]: Leaving directory '/home/wyj/Code/C/GR_4d_raytracing'
OMP_NUM_THREADS=16 build/Release/schwarzschild_sky --catalog assets/sky_grid_5deg.csv --output /tmp/gr_psf_event_sink_reference/schwarzschild_ra1_dec1_fov60_640x360.png --hdr-output --width 640 --height 360 --fov-deg 60 --look-ra-deg 1 --look-dec-deg 1 --exposure 0.1 --observer-radius 30 --observer-velocity 0 0 0 --psf-fwhm-pixels 2.7 --psf-moffat-beta 4.5 --max-magnification 1e300 --max-cache-psf-flux 1 --psf-relative-tail 1e-8 --psf-min-y 0 --coarse-cell-pixels 16 --refine-max-level 0 --refine-angle-abs-deg 0.001 --refine-angle-rel 0.1 --refine-jacobian-min 1e-3 --refine-min-edge-pixels 0.5 --refine-min-area-pixels2 0.25 --catalog-load-workers 4
PSF cache ready: 64x64 phases, radius 47 px, relative tail 1e-08, tail abs 1e-06, boundary 1e-07, build 0.670 s
Rendered 5206 images from 5654 catalog stars to /tmp/gr_psf_event_sink_reference/schwarzschild_ra1_dec1_fov60_640x360.png (ok)
PSF splats: cached 5206, cached wing-clipped 0, direct fallbacks 0, discarded below min-Y 0
python3 scripts/fits_floatdiff.py tests/data/psf_event_sink_reference/schwarzschild_ra1_dec1_fov60_640x360_HDR.fits /tmp/gr_psf_event_sink_reference/schwarzschild_ra1_dec1_fov60_640x360_HDR.fits
floatdiff: shape=(640, 360, 3) samples=691200 mismatches=0 max_abs=0 max_rel=0 worst_sample=0
./build/Release/test_observer_minkowski
coordinate-camera regression passed
./build/Release/test_observer_schwarzschild
coordinate-camera regression passed
./build/Release/test_geodesic
./build/Release/test_frame
./build/Release/test_schwarzschild
./build/Release/test_observer_track
./build/Release/test_catalog_prefetch
python3 tests/test_camera_cli.py build/Release
minkowski: CLI checks passed; single/movie PNG identical, map max error 0
schwarzschild: movie lens map identical with 1, 4 and 16 threads
schwarzschild: CLI checks passed; single/movie PNG identical, map max error 4.44e-16
benchmarks/hip_psf_2026-09-07/fail_fast.log
+3
View File
@@ -0,0 +1,3 @@
PSF cache ready: 64x64 phases, radius 47 px, relative tail 1e-08, tail abs 1e-06, boundary 1e-07, build 1.636 s
HIP PSF backend failed during initialization: no ROCm-capable device is detected
exit_code=1
benchmarks/hip_psf_2026-09-07/final_indexed.log
+208
View File
@@ -0,0 +1,208 @@
PSF cache ready: 64x64 phases, radius 47 px, relative tail 1e-08, tail abs 1e-06, boundary 1e-07, build 0.674 s
Frame 0: finding and prefetching catalog tiles...
Catalog prefetch: 64440 requested, 1 newly loaded (8192 stars), 64439 unavailable in 0.034 s; 4 loader workers
Frame 0: catalog prefetch finished (64440 candidate tiles).
Frame 0: splatting 130882 lens triangles...
Frame 0: splat worker 13/16 started.
Frame 0: splat worker 1/16 started.
Frame 0: splat worker 9/16 started.
Frame 0: splat worker 13/16 reached 8 local triangles.
Frame 0: splat worker 4/16 started.
Frame 0: splat worker 16/16 started.
Frame 0: splat worker 15/16 started.
Frame 0: splat worker 16/16 reached 8 local triangles.
Frame 0: splat worker 11/16 started.
Frame 0: splat worker 3/16 started.
Frame 0: splat worker 7/16 started.
Frame 0: splat worker 1/16 reached 8 local triangles.
Frame 0: splat worker 3/16 reached 8 local triangles.
Frame 0: splat worker 1/16 reached 16 local triangles.
Frame 0: splat worker 12/16 started.
Frame 0: splat worker 5/16 started.
Frame 0: splat worker 9/16 reached 8 local triangles.
Frame 0: splat worker 12/16 reached 8 local triangles.
Frame 0: splat worker 9/16 reached 16 local triangles.
Frame 0: splat worker 12/16 reached 16 local triangles.
Frame 0: splat worker 4/16 reached 8 local triangles.
Frame 0: splat worker 2/16 started.
Frame 0: splat worker 16/16 reached 16 local triangles.
Frame 0: splat worker 15/16 reached 8 local triangles.
Frame 0: splat worker 11/16 reached 8 local triangles.
Frame 0: splat worker 15/16 reached 16 local triangles.
Frame 0: splat worker 8/16 started.
Frame 0: splat worker 11/16 reached 16 local triangles.
Frame 0: splat worker 7/16 reached 8 local triangles.
Frame 0: splat worker 15/16 reached 32 local triangles.
Frame 0: splat worker 8/16 reached 8 local triangles.
Frame 0: splat worker 7/16 reached 16 local triangles.
Frame 0: splat worker 13/16 reached 16 local triangles.
Frame 0: splat worker 10/16 started.
Frame 0: splat worker 4/16 reached 16 local triangles.
Frame 0: splat worker 2/16 reached 8 local triangles.
Frame 0: splat worker 10/16 reached 8 local triangles.
Frame 0: splat worker 2/16 reached 16 local triangles.
Frame 0: splat worker 10/16 reached 16 local triangles.
Frame 0: splat worker 16/16 reached 32 local triangles.
Frame 0: splat worker 3/16 reached 16 local triangles.
Frame 0: splat worker 1/16 reached 32 local triangles.
Frame 0: splat worker 11/16 reached 32 local triangles.
Frame 0: splat worker 5/16 reached 8 local triangles.
Frame 0: splat worker 8/16 reached 16 local triangles.
Frame 0: splat worker 5/16 reached 16 local triangles.
Frame 0: splat worker 1/16 reached 64 local triangles.
Frame 0: splat worker 11/16 reached 64 local triangles.
Frame 0: splat worker 8/16 reached 32 local triangles.
Frame 0: splat worker 9/16 reached 32 local triangles.
Frame 0: splat worker 12/16 reached 32 local triangles.
Frame 0: splat worker 4/16 reached 32 local triangles.
Frame 0: splat worker 6/16 started.
Frame 0: splat worker 2/16 reached 32 local triangles.
Frame 0: splat worker 1/16 reached 128 local triangles.
Frame 0: splat worker 9/16 reached 64 local triangles.
Frame 0: splat worker 12/16 reached 64 local triangles.
Frame 0: splat worker 14/16 started.
Frame 0: splat worker 16/16 reached 64 local triangles.
Frame 0: splat worker 7/16 reached 32 local triangles.
Frame 0: splat worker 13/16 reached 32 local triangles.
Frame 0: splat worker 10/16 reached 32 local triangles.
Frame 0: splat worker 7/16 reached 64 local triangles.
Frame 0: splat worker 6/16 reached 8 local triangles.
Frame 0: splat worker 3/16 reached 32 local triangles.
Frame 0: splat worker 6/16 reached 16 local triangles.
Frame 0: splat worker 4/16 reached 64 local triangles.
Frame 0: splat worker 13/16 reached 64 local triangles.
Frame 0: splat worker 10/16 reached 64 local triangles.
Frame 0: splat worker 14/16 reached 8 local triangles.
Frame 0: splat worker 15/16 reached 64 local triangles.
Frame 0: splat worker 14/16 reached 16 local triangles.
Frame 0: splat worker 1/16 reached 256 local triangles.
Frame 0: splat worker 16/16 reached 128 local triangles.
Frame 0: splat worker 9/16 reached 128 local triangles.
Frame 0: splat worker 2/16 reached 64 local triangles.
Frame 0: splat worker 12/16 reached 128 local triangles.
Frame 0: splat worker 6/16 reached 32 local triangles.
Frame 0: splat worker 5/16 reached 32 local triangles.
Frame 0: splat worker 3/16 reached 64 local triangles.
Frame 0: splat worker 15/16 reached 128 local triangles.
Frame 0: splat worker 11/16 reached 128 local triangles.
Frame 0: splat worker 7/16 reached 128 local triangles.
Frame 0: splat worker 14/16 reached 32 local triangles.
Frame 0: splat worker 4/16 reached 128 local triangles.
Frame 0: splat worker 13/16 reached 128 local triangles.
Frame 0: splat worker 10/16 reached 128 local triangles.
Frame 0: splat worker 8/16 reached 64 local triangles.
Frame 0: splat worker 2/16 reached 128 local triangles.
Frame 0: splat worker 14/16 reached 64 local triangles.
Frame 0: splat worker 16/16 reached 256 local triangles.
Frame 0: splat worker 3/16 reached 128 local triangles.
Frame 0: splat worker 5/16 reached 64 local triangles.
Frame 0: splat worker 9/16 reached 256 local triangles.
Frame 0: splat worker 8/16 reached 128 local triangles.
Frame 0: splat worker 12/16 reached 256 local triangles.
Frame 0: splat worker 15/16 reached 256 local triangles.
Frame 0: splat worker 14/16 reached 128 local triangles.
Frame 0: splat worker 11/16 reached 256 local triangles.
Frame 0: splat worker 7/16 reached 256 local triangles.
Frame 0: splat worker 2/16 reached 256 local triangles.
Frame 0: splat worker 5/16 reached 128 local triangles.
Frame 0: splat worker 13/16 reached 256 local triangles.
Frame 0: splat worker 4/16 reached 256 local triangles.
Frame 0: splat worker 10/16 reached 256 local triangles.
Frame 0: splat worker 1/16 reached 512 local triangles.
Frame 0: splat worker 3/16 reached 256 local triangles.
Frame 0: splat worker 8/16 reached 256 local triangles.
Frame 0: splat worker 14/16 reached 256 local triangles.
Frame 0: splat worker 5/16 reached 256 local triangles.
Frame 0: splat worker 16/16 reached 512 local triangles.
Frame 0: splat worker 9/16 reached 512 local triangles.
Frame 0: splat worker 12/16 reached 512 local triangles.
Frame 0: splat worker 15/16 reached 512 local triangles.
Frame 0: splat worker 7/16 reached 512 local triangles.
Frame 0: splat worker 13/16 reached 512 local triangles.
Frame 0: splat worker 11/16 reached 512 local triangles.
Frame 0: splat worker 2/16 reached 512 local triangles.
Frame 0: splat worker 4/16 reached 512 local triangles.
Frame 0: splat worker 10/16 reached 512 local triangles.
Frame 0: splat worker 3/16 reached 512 local triangles.
Frame 0: splat worker 8/16 reached 512 local triangles.
Frame 0: splat worker 14/16 reached 512 local triangles.
Frame 0: splat worker 5/16 reached 512 local triangles.
Frame 0: splat worker 1/16 reached 1024 local triangles.
Frame 0: splat worker 13/16 reached 1024 local triangles.
Frame 0: splat worker 16/16 reached 1024 local triangles.
Frame 0: splat worker 9/16 reached 1024 local triangles.
Frame 0: splat worker 12/16 reached 1024 local triangles.
Frame 0: splat worker 7/16 reached 1024 local triangles.
Frame 0: splat worker 15/16 reached 1024 local triangles.
Frame 0: splat worker 11/16 reached 1024 local triangles.
Frame 0: splat worker 2/16 reached 1024 local triangles.
Frame 0: splat worker 4/16 reached 1024 local triangles.
Frame 0: splat worker 10/16 reached 1024 local triangles.
Frame 0: splat worker 8/16 reached 1024 local triangles.
Frame 0: splat worker 14/16 reached 1024 local triangles.
Frame 0: splat worker 3/16 reached 1024 local triangles.
Frame 0: splat worker 5/16 reached 1024 local triangles.
Frame 0: splat worker 16/16 reached 2048 local triangles.
Frame 0: splat worker 11/16 reached 2048 local triangles.
Frame 0: splat worker 2/16 reached 2048 local triangles.
Frame 0: splat worker 9/16 reached 2048 local triangles.
Frame 0: splat worker 13/16 reached 2048 local triangles.
Frame 0: splat worker 12/16 reached 2048 local triangles.
Frame 0: splat worker 16/16 reached 4096 local triangles.
Frame 0: splat worker 2/16 reached 4096 local triangles.
Frame 0: splat worker 11/16 reached 4096 local triangles.
Frame 0: splat worker 1/16 reached 2048 local triangles.
Frame 0: splat worker 6/16 reached 64 local triangles.
Frame 0: splat worker 9/16 reached 4096 local triangles.
Frame 0: splat worker 13/16 reached 4096 local triangles.
Frame 0: splat worker 6/16 reached 128 local triangles.
Frame 0: splat worker 6/16 reached 256 local triangles.
Frame 0: splat worker 5/16 reached 2048 local triangles.
Frame 0: splat worker 6/16 reached 512 local triangles.
Frame 0: splat worker 8/16 reached 2048 local triangles.
Frame 0: splat worker 6/16 reached 1024 local triangles.
Frame 0: splat worker 4/16 reached 2048 local triangles.
Frame 0: splat worker 12/16 reached 4096 local triangles.
Frame 0: splat worker 10/16 reached 2048 local triangles.
Frame 0: splat worker 3/16 reached 2048 local triangles.
Frame 0: splat worker 1/16 reached 4096 local triangles.
Frame 0: splat worker 14/16 reached 2048 local triangles.
Frame 0: splat worker 7/16 reached 2048 local triangles.
Frame 0: splat worker 4/16 reached 4096 local triangles.
Frame 0: splat worker 8/16 reached 4096 local triangles.
Frame 0: splat worker 5/16 reached 4096 local triangles.
Frame 0: splat worker 13/16 reached 8192 local triangles.
Frame 0: splat worker 10/16 reached 4096 local triangles.
Frame 0: splat worker 16/16 reached 8192 local triangles.
Frame 0: splat worker 3/16 reached 4096 local triangles.
Frame 0: splat worker 11/16 reached 8192 local triangles.
Frame 0: splat worker 6/16 reached 2048 local triangles.
Frame 0: splat worker 14/16 reached 4096 local triangles.
Frame 0: splat worker 9/16 reached 8192 local triangles.
Frame 0: splat worker 7/16 reached 4096 local triangles.
Frame 0: splat worker 15/16 reached 2048 local triangles.
Frame 0: splat worker 12/16 reached 8192 local triangles.
Frame 0: splat worker 1/16 reached 8192 local triangles.
Frame 0: splat worker 8/16 reached 8192 local triangles.
Frame 0: splat worker 10/16 finished after 7675 local triangles.
Frame 0: splat worker 15/16 finished after 3525 local triangles.
Frame 0: splat worker 8/16 finished after 8466 local triangles.
Frame 0: splat worker 14/16 finished after 6755 local triangles.
Frame 0: splat worker 3/16 finished after 7436 local triangles.
Frame 0: splat worker 1/16 finished after 9593 local triangles.
Frame 0: splat worker 5/16 finished after 8031 local triangles.
Frame 0: splat worker 11/16 finished after 11275 local triangles.
Frame 0: splat worker 2/16 finished after 8050 local triangles.
Frame 0: splat worker 13/16 finished after 13527 local triangles.
Frame 0: splat worker 9/16 finished after 10796 local triangles.
Frame 0: splat worker 12/16 finished after 9604 local triangles.
Frame 0: splat worker 7/16 finished after 6246 local triangles.
Frame 0: splat worker 4/16 finished after 8025 local triangles.
Frame 0: splat worker 6/16 finished after 3431 local triangles.
Frame 0: splat worker 16/16 finished after 8447 local triangles.
HIP producers: 16 workers; summed generation 0.218 s, submission/fallback 0.205 s; wall 0.066 s
Frame 0: catalog splatting finished in 0.1 s; writing image...
Rendered 12765 images from 8192 catalog stars to /tmp/gr-hip-final-indexed.png (ok; imported lens map)
PSF splats: cached 12765, cached wing-clipped 0, direct fallbacks 0, discarded below min-Y 0
HIP PSF: 12765 events in 15 batches (15 timed); upload 0.000218 s, kernel 0.028882 s, download 0.023217 s
OMP_NUM_THREADS=16 timeout --kill-after=5s 60s --all-sky-catalog 4 1e-8 11.10s user 0.26s system 940% cpu 1.208 total
benchmarks/hip_psf_2026-09-07/final_lens8192.log
+210
View File
@@ -0,0 +1,210 @@
PSF cache ready: 64x64 phases, radius 47 px, relative tail 1e-08, tail abs 1e-06, boundary 1e-07, build 0.728 s
Frame 0: finding and prefetching catalog tiles...
Frame 0: catalog prefetch finished (0 candidate tiles).
Frame 0: splatting 130882 lens triangles...
Frame 0: splat worker 1/16 started.
Frame 0: splat worker 3/16 started.
Frame 0: splat worker 4/16 started.
Frame 0: splat worker 9/16 started.
Frame 0: splat worker 13/16 started.
Frame 0: splat worker 2/16 started.
Frame 0: splat worker 7/16 started.
Frame 0: splat worker 14/16 started.
Frame 0: splat worker 10/16 started.
Frame 0: splat worker 15/16 started.
Frame 0: splat worker 8/16 started.
Frame 0: splat worker 12/16 started.
Frame 0: splat worker 5/16 started.
Frame 0: splat worker 11/16 started.
Frame 0: splat worker 7/16 reached 8 local triangles.
Frame 0: splat worker 1/16 reached 8 local triangles.
Frame 0: splat worker 13/16 reached 8 local triangles.
Frame 0: splat worker 3/16 reached 8 local triangles.
Frame 0: splat worker 14/16 reached 8 local triangles.
Frame 0: splat worker 8/16 reached 8 local triangles.
Frame 0: splat worker 4/16 reached 8 local triangles.
Frame 0: splat worker 15/16 reached 8 local triangles.
Frame 0: splat worker 9/16 reached 8 local triangles.
Frame 0: splat worker 2/16 reached 8 local triangles.
Frame 0: splat worker 10/16 reached 8 local triangles.
Frame 0: splat worker 11/16 reached 8 local triangles.
Frame 0: splat worker 16/16 started.
Frame 0: splat worker 6/16 started.
Frame 0: splat worker 12/16 reached 8 local triangles.
Frame 0: splat worker 5/16 reached 8 local triangles.
Frame 0: splat worker 7/16 reached 16 local triangles.
Frame 0: splat worker 6/16 reached 8 local triangles.
Frame 0: splat worker 13/16 reached 16 local triangles.
Frame 0: splat worker 10/16 reached 16 local triangles.
Frame 0: splat worker 4/16 reached 16 local triangles.
Frame 0: splat worker 14/16 reached 16 local triangles.
Frame 0: splat worker 1/16 reached 16 local triangles.
Frame 0: splat worker 3/16 reached 16 local triangles.
Frame 0: splat worker 15/16 reached 16 local triangles.
Frame 0: splat worker 8/16 reached 16 local triangles.
Frame 0: splat worker 9/16 reached 16 local triangles.
Frame 0: splat worker 2/16 reached 16 local triangles.
Frame 0: splat worker 12/16 reached 16 local triangles.
Frame 0: splat worker 5/16 reached 16 local triangles.
Frame 0: splat worker 16/16 reached 8 local triangles.
Frame 0: splat worker 6/16 reached 16 local triangles.
Frame 0: splat worker 16/16 reached 16 local triangles.
Frame 0: splat worker 7/16 reached 32 local triangles.
Frame 0: splat worker 6/16 reached 32 local triangles.
Frame 0: splat worker 4/16 reached 32 local triangles.
Frame 0: splat worker 13/16 reached 32 local triangles.
Frame 0: splat worker 3/16 reached 32 local triangles.
Frame 0: splat worker 10/16 reached 32 local triangles.
Frame 0: splat worker 14/16 reached 32 local triangles.
Frame 0: splat worker 12/16 reached 32 local triangles.
Frame 0: splat worker 9/16 reached 32 local triangles.
Frame 0: splat worker 1/16 reached 32 local triangles.
Frame 0: splat worker 5/16 reached 32 local triangles.
Frame 0: splat worker 8/16 reached 32 local triangles.
Frame 0: splat worker 15/16 reached 32 local triangles.
Frame 0: splat worker 2/16 reached 32 local triangles.
Frame 0: splat worker 16/16 reached 32 local triangles.
Frame 0: splat worker 6/16 reached 64 local triangles.
Frame 0: splat worker 11/16 reached 16 local triangles.
Frame 0: splat worker 10/16 reached 64 local triangles.
Frame 0: splat worker 4/16 reached 64 local triangles.
Frame 0: splat worker 8/16 reached 64 local triangles.
Frame 0: splat worker 13/16 reached 64 local triangles.
Frame 0: splat worker 7/16 reached 64 local triangles.
Frame 0: splat worker 14/16 reached 64 local triangles.
Frame 0: splat worker 2/16 reached 64 local triangles.
Frame 0: splat worker 12/16 reached 64 local triangles.
Frame 0: splat worker 9/16 reached 64 local triangles.
Frame 0: splat worker 3/16 reached 64 local triangles.
Frame 0: splat worker 5/16 reached 64 local triangles.
Frame 0: splat worker 15/16 reached 64 local triangles.
Frame 0: splat worker 1/16 reached 64 local triangles.
Frame 0: splat worker 11/16 reached 32 local triangles.
Frame 0: splat worker 6/16 reached 128 local triangles.
Frame 0: splat worker 16/16 reached 64 local triangles.
Frame 0: splat worker 8/16 reached 128 local triangles.
Frame 0: splat worker 2/16 reached 128 local triangles.
Frame 0: splat worker 10/16 reached 128 local triangles.
Frame 0: splat worker 13/16 reached 128 local triangles.
Frame 0: splat worker 15/16 reached 128 local triangles.
Frame 0: splat worker 4/16 reached 128 local triangles.
Frame 0: splat worker 14/16 reached 128 local triangles.
Frame 0: splat worker 9/16 reached 128 local triangles.
Frame 0: splat worker 5/16 reached 128 local triangles.
Frame 0: splat worker 7/16 reached 128 local triangles.
Frame 0: splat worker 1/16 reached 128 local triangles.
Frame 0: splat worker 12/16 reached 128 local triangles.
Frame 0: splat worker 3/16 reached 128 local triangles.
Frame 0: splat worker 11/16 reached 64 local triangles.
Frame 0: splat worker 6/16 reached 256 local triangles.
Frame 0: splat worker 16/16 reached 128 local triangles.
Frame 0: splat worker 8/16 reached 256 local triangles.
Frame 0: splat worker 13/16 reached 256 local triangles.
Frame 0: splat worker 15/16 reached 256 local triangles.
Frame 0: splat worker 5/16 reached 256 local triangles.
Frame 0: splat worker 4/16 reached 256 local triangles.
Frame 0: splat worker 1/16 reached 256 local triangles.
Frame 0: splat worker 2/16 reached 256 local triangles.
Frame 0: splat worker 14/16 reached 256 local triangles.
Frame 0: splat worker 7/16 reached 256 local triangles.
Frame 0: splat worker 12/16 reached 256 local triangles.
Frame 0: splat worker 10/16 reached 256 local triangles.
Frame 0: splat worker 3/16 reached 256 local triangles.
Frame 0: splat worker 11/16 reached 128 local triangles.
Frame 0: splat worker 9/16 reached 256 local triangles.
Frame 0: splat worker 6/16 reached 512 local triangles.
Frame 0: splat worker 15/16 reached 512 local triangles.
Frame 0: splat worker 8/16 reached 512 local triangles.
Frame 0: splat worker 13/16 reached 512 local triangles.
Frame 0: splat worker 4/16 reached 512 local triangles.
Frame 0: splat worker 12/16 reached 512 local triangles.
Frame 0: splat worker 5/16 reached 512 local triangles.
Frame 0: splat worker 16/16 reached 256 local triangles.
Frame 0: splat worker 1/16 reached 512 local triangles.
Frame 0: splat worker 2/16 reached 512 local triangles.
Frame 0: splat worker 10/16 reached 512 local triangles.
Frame 0: splat worker 14/16 reached 512 local triangles.
Frame 0: splat worker 3/16 reached 512 local triangles.
Frame 0: splat worker 9/16 reached 512 local triangles.
Frame 0: splat worker 7/16 reached 512 local triangles.
Frame 0: splat worker 11/16 reached 256 local triangles.
Frame 0: splat worker 6/16 reached 1024 local triangles.
Frame 0: splat worker 16/16 reached 512 local triangles.
Frame 0: splat worker 11/16 reached 512 local triangles.
Frame 0: splat worker 13/16 reached 1024 local triangles.
Frame 0: splat worker 15/16 reached 1024 local triangles.
Frame 0: splat worker 1/16 reached 1024 local triangles.
Frame 0: splat worker 8/16 reached 1024 local triangles.
Frame 0: splat worker 4/16 reached 1024 local triangles.
Frame 0: splat worker 5/16 reached 1024 local triangles.
Frame 0: splat worker 12/16 reached 1024 local triangles.
Frame 0: splat worker 10/16 reached 1024 local triangles.
Frame 0: splat worker 14/16 reached 1024 local triangles.
Frame 0: splat worker 2/16 reached 1024 local triangles.
Frame 0: splat worker 3/16 reached 1024 local triangles.
Frame 0: splat worker 9/16 reached 1024 local triangles.
Frame 0: splat worker 7/16 reached 1024 local triangles.
Frame 0: splat worker 16/16 reached 1024 local triangles.
Frame 0: splat worker 11/16 reached 1024 local triangles.
Frame 0: splat worker 6/16 reached 2048 local triangles.
Frame 0: splat worker 15/16 reached 2048 local triangles.
Frame 0: splat worker 13/16 reached 2048 local triangles.
Frame 0: splat worker 12/16 reached 2048 local triangles.
Frame 0: splat worker 5/16 reached 2048 local triangles.
Frame 0: splat worker 14/16 reached 2048 local triangles.
Frame 0: splat worker 1/16 reached 2048 local triangles.
Frame 0: splat worker 2/16 reached 2048 local triangles.
Frame 0: splat worker 3/16 reached 2048 local triangles.
Frame 0: splat worker 4/16 reached 2048 local triangles.
Frame 0: splat worker 9/16 reached 2048 local triangles.
Frame 0: splat worker 7/16 reached 2048 local triangles.
Frame 0: splat worker 8/16 reached 2048 local triangles.
Frame 0: splat worker 10/16 reached 2048 local triangles.
Frame 0: splat worker 16/16 reached 2048 local triangles.
Frame 0: splat worker 11/16 reached 2048 local triangles.
Frame 0: splat worker 6/16 reached 4096 local triangles.
Frame 0: splat worker 13/16 reached 4096 local triangles.
Frame 0: splat worker 15/16 reached 4096 local triangles.
Frame 0: splat worker 8/16 reached 4096 local triangles.
Frame 0: splat worker 9/16 reached 4096 local triangles.
Frame 0: splat worker 12/16 reached 4096 local triangles.
Frame 0: splat worker 5/16 reached 4096 local triangles.
Frame 0: splat worker 4/16 reached 4096 local triangles.
Frame 0: splat worker 2/16 reached 4096 local triangles.
Frame 0: splat worker 1/16 reached 4096 local triangles.
Frame 0: splat worker 10/16 reached 4096 local triangles.
Frame 0: splat worker 3/16 reached 4096 local triangles.
Frame 0: splat worker 14/16 reached 4096 local triangles.
Frame 0: splat worker 7/16 reached 4096 local triangles.
Frame 0: splat worker 16/16 reached 4096 local triangles.
Frame 0: splat worker 11/16 reached 4096 local triangles.
Frame 0: splat worker 5/16 reached 8192 local triangles.
Frame 0: splat worker 6/16 reached 8192 local triangles.
Frame 0: splat worker 15/16 reached 8192 local triangles.
Frame 0: splat worker 13/16 reached 8192 local triangles.
Frame 0: splat worker 1/16 reached 8192 local triangles.
Frame 0: splat worker 12/16 reached 8192 local triangles.
Frame 0: splat worker 9/16 reached 8192 local triangles.
Frame 0: splat worker 10/16 reached 8192 local triangles.
Frame 0: splat worker 3/16 reached 8192 local triangles.
Frame 0: splat worker 1/16 finished after 8417 local triangles.
Frame 0: splat worker 3/16 finished after 8220 local triangles.
Frame 0: splat worker 13/16 finished after 8435 local triangles.
Frame 0: splat worker 5/16 finished after 8644 local triangles.
Frame 0: splat worker 9/16 finished after 8288 local triangles.
Frame 0: splat worker 8/16 finished after 8175 local triangles.
Frame 0: splat worker 14/16 finished after 7891 local triangles.
Frame 0: splat worker 6/16 finished after 8634 local triangles.
Frame 0: splat worker 11/16 finished after 7108 local triangles.
Frame 0: splat worker 7/16 finished after 7829 local triangles.
Frame 0: splat worker 10/16 finished after 8276 local triangles.
Frame 0: splat worker 4/16 finished after 8133 local triangles.
Frame 0: splat worker 15/16 finished after 8542 local triangles.
Frame 0: splat worker 16/16 finished after 7993 local triangles.
Frame 0: splat worker 2/16 finished after 7920 local triangles.
Frame 0: splat worker 12/16 finished after 8377 local triangles.
Frame 0: catalog splatting finished in 0.6 s; writing image...
Rendered 12765 images from 8192 catalog stars to /tmp/gr-hip-final-8192.png (ok; imported lens map)
PSF splats: cached 12765, cached wing-clipped 0, direct fallbacks 0, discarded below min-Y 0
HIP PSF: 12765 events in 14 batches (14 timed); upload 0.000188 s, kernel 0.023472 s, download 0.025256 s
OMP_NUM_THREADS=16 timeout --kill-after=5s 60s --catalog --psf-relative-tai 18.45s user 0.19s system 1092% cpu 1.706 total
benchmarks/hip_psf_2026-09-07/final_multi.log
+210
View File
@@ -0,0 +1,210 @@
PSF cache ready: 64x64 phases, radius 47 px, relative tail 1e-08, tail abs 1e-06, boundary 1e-07, build 0.674 s
Frame 0: finding and prefetching catalog tiles...
Catalog prefetch: 64440 requested, 3 newly loaded (21087 stars), 64437 unavailable in 0.033 s; 4 loader workers
Frame 0: catalog prefetch finished (64440 candidate tiles).
Frame 0: splatting 130882 lens triangles...
Frame 0: splat worker 1/16 started.
Frame 0: splat worker 9/16 started.
Frame 0: splat worker 12/16 started.
Frame 0: splat worker 13/16 started.
Frame 0: splat worker 8/16 started.
Frame 0: splat worker 8/16 reached 8 local triangles.
Frame 0: splat worker 8/16 reached 16 local triangles.
Frame 0: splat worker 6/16 started.
Frame 0: splat worker 6/16 reached 8 local triangles.
Frame 0: splat worker 6/16 reached 16 local triangles.
Frame 0: splat worker 6/16 reached 32 local triangles.
Frame 0: splat worker 6/16 reached 64 local triangles.
Frame 0: splat worker 5/16 started.
Frame 0: splat worker 5/16 reached 8 local triangles.
Frame 0: splat worker 5/16 reached 16 local triangles.
Frame 0: splat worker 5/16 reached 32 local triangles.
Frame 0: splat worker 5/16 reached 64 local triangles.
Frame 0: splat worker 5/16 reached 128 local triangles.
Frame 0: splat worker 7/16 started.
Frame 0: splat worker 1/16 reached 8 local triangles.
Frame 0: splat worker 7/16 reached 8 local triangles.
Frame 0: splat worker 16/16 started.
Frame 0: splat worker 9/16 reached 8 local triangles.
Frame 0: splat worker 15/16 started.
Frame 0: splat worker 9/16 reached 16 local triangles.
Frame 0: splat worker 14/16 started.
Frame 0: splat worker 9/16 reached 32 local triangles.
Frame 0: splat worker 14/16 reached 8 local triangles.
Frame 0: splat worker 4/16 started.
Frame 0: splat worker 9/16 reached 64 local triangles.
Frame 0: splat worker 13/16 reached 8 local triangles.
Frame 0: splat worker 2/16 started.
Frame 0: splat worker 13/16 reached 16 local triangles.
Frame 0: splat worker 2/16 reached 8 local triangles.
Frame 0: splat worker 9/16 reached 128 local triangles.
Frame 0: splat worker 2/16 reached 16 local triangles.
Frame 0: splat worker 10/16 started.
Frame 0: splat worker 1/16 reached 16 local triangles.
Frame 0: splat worker 7/16 reached 16 local triangles.
Frame 0: splat worker 3/16 started.
Frame 0: splat worker 16/16 reached 8 local triangles.
Frame 0: splat worker 11/16 started.
Frame 0: splat worker 5/16 reached 256 local triangles.
Frame 0: splat worker 14/16 reached 16 local triangles.
Frame 0: splat worker 4/16 reached 8 local triangles.
Frame 0: splat worker 15/16 reached 8 local triangles.
Frame 0: splat worker 13/16 reached 32 local triangles.
Frame 0: splat worker 6/16 reached 128 local triangles.
Frame 0: splat worker 12/16 reached 8 local triangles.
Frame 0: splat worker 16/16 reached 16 local triangles.
Frame 0: splat worker 16/16 reached 32 local triangles.
Frame 0: splat worker 11/16 reached 8 local triangles.
Frame 0: splat worker 16/16 reached 64 local triangles.
Frame 0: splat worker 4/16 reached 16 local triangles.
Frame 0: splat worker 11/16 reached 16 local triangles.
Frame 0: splat worker 4/16 reached 32 local triangles.
Frame 0: splat worker 11/16 reached 32 local triangles.
Frame 0: splat worker 14/16 reached 32 local triangles.
Frame 0: splat worker 12/16 reached 16 local triangles.
Frame 0: splat worker 10/16 reached 8 local triangles.
Frame 0: splat worker 1/16 reached 32 local triangles.
Frame 0: splat worker 4/16 reached 64 local triangles.
Frame 0: splat worker 2/16 reached 32 local triangles.
Frame 0: splat worker 16/16 reached 128 local triangles.
Frame 0: splat worker 3/16 reached 8 local triangles.
Frame 0: splat worker 7/16 reached 32 local triangles.
Frame 0: splat worker 11/16 reached 64 local triangles.
Frame 0: splat worker 14/16 reached 64 local triangles.
Frame 0: splat worker 1/16 reached 64 local triangles.
Frame 0: splat worker 15/16 reached 16 local triangles.
Frame 0: splat worker 13/16 reached 64 local triangles.
Frame 0: splat worker 12/16 reached 32 local triangles.
Frame 0: splat worker 14/16 reached 128 local triangles.
Frame 0: splat worker 4/16 reached 128 local triangles.
Frame 0: splat worker 7/16 reached 64 local triangles.
Frame 0: splat worker 15/16 reached 32 local triangles.
Frame 0: splat worker 3/16 reached 16 local triangles.
Frame 0: splat worker 3/16 reached 32 local triangles.
Frame 0: splat worker 1/16 reached 128 local triangles.
Frame 0: splat worker 3/16 reached 64 local triangles.
Frame 0: splat worker 12/16 reached 64 local triangles.
Frame 0: splat worker 9/16 reached 256 local triangles.
Frame 0: splat worker 10/16 reached 16 local triangles.
Frame 0: splat worker 15/16 reached 64 local triangles.
Frame 0: splat worker 10/16 reached 32 local triangles.
Frame 0: splat worker 3/16 reached 128 local triangles.
Frame 0: splat worker 8/16 reached 32 local triangles.
Frame 0: splat worker 2/16 reached 64 local triangles.
Frame 0: splat worker 11/16 reached 128 local triangles.
Frame 0: splat worker 10/16 reached 64 local triangles.
Frame 0: splat worker 1/16 reached 256 local triangles.
Frame 0: splat worker 12/16 reached 128 local triangles.
Frame 0: splat worker 4/16 reached 256 local triangles.
Frame 0: splat worker 8/16 reached 64 local triangles.
Frame 0: splat worker 7/16 reached 128 local triangles.
Frame 0: splat worker 16/16 reached 256 local triangles.
Frame 0: splat worker 13/16 reached 128 local triangles.
Frame 0: splat worker 3/16 reached 256 local triangles.
Frame 0: splat worker 6/16 reached 256 local triangles.
Frame 0: splat worker 13/16 reached 256 local triangles.
Frame 0: splat worker 15/16 reached 128 local triangles.
Frame 0: splat worker 2/16 reached 128 local triangles.
Frame 0: splat worker 11/16 reached 256 local triangles.
Frame 0: splat worker 7/16 reached 256 local triangles.
Frame 0: splat worker 9/16 reached 512 local triangles.
Frame 0: splat worker 5/16 reached 512 local triangles.
Frame 0: splat worker 1/16 reached 512 local triangles.
Frame 0: splat worker 10/16 reached 128 local triangles.
Frame 0: splat worker 8/16 reached 128 local triangles.
Frame 0: splat worker 14/16 reached 256 local triangles.
Frame 0: splat worker 3/16 reached 512 local triangles.
Frame 0: splat worker 10/16 reached 256 local triangles.
Frame 0: splat worker 15/16 reached 256 local triangles.
Frame 0: splat worker 12/16 reached 256 local triangles.
Frame 0: splat worker 4/16 reached 512 local triangles.
Frame 0: splat worker 2/16 reached 256 local triangles.
Frame 0: splat worker 6/16 reached 512 local triangles.
Frame 0: splat worker 11/16 reached 512 local triangles.
Frame 0: splat worker 16/16 reached 512 local triangles.
Frame 0: splat worker 14/16 reached 512 local triangles.
Frame 0: splat worker 12/16 reached 512 local triangles.
Frame 0: splat worker 8/16 reached 256 local triangles.
Frame 0: splat worker 10/16 reached 512 local triangles.
Frame 0: splat worker 3/16 reached 1024 local triangles.
Frame 0: splat worker 15/16 reached 512 local triangles.
Frame 0: splat worker 7/16 reached 512 local triangles.
Frame 0: splat worker 13/16 reached 512 local triangles.
Frame 0: splat worker 5/16 reached 1024 local triangles.
Frame 0: splat worker 9/16 reached 1024 local triangles.
Frame 0: splat worker 4/16 reached 1024 local triangles.
Frame 0: splat worker 1/16 reached 1024 local triangles.
Frame 0: splat worker 2/16 reached 512 local triangles.
Frame 0: splat worker 16/16 reached 1024 local triangles.
Frame 0: splat worker 6/16 reached 1024 local triangles.
Frame 0: splat worker 11/16 reached 1024 local triangles.
Frame 0: splat worker 15/16 reached 1024 local triangles.
Frame 0: splat worker 8/16 reached 512 local triangles.
Frame 0: splat worker 14/16 reached 1024 local triangles.
Frame 0: splat worker 12/16 reached 1024 local triangles.
Frame 0: splat worker 3/16 reached 2048 local triangles.
Frame 0: splat worker 10/16 reached 1024 local triangles.
Frame 0: splat worker 7/16 reached 1024 local triangles.
Frame 0: splat worker 2/16 reached 1024 local triangles.
Frame 0: splat worker 5/16 reached 2048 local triangles.
Frame 0: splat worker 16/16 reached 2048 local triangles.
Frame 0: splat worker 4/16 reached 2048 local triangles.
Frame 0: splat worker 11/16 reached 2048 local triangles.
Frame 0: splat worker 6/16 reached 2048 local triangles.
Frame 0: splat worker 14/16 reached 2048 local triangles.
Frame 0: splat worker 3/16 reached 4096 local triangles.
Frame 0: splat worker 12/16 reached 2048 local triangles.
Frame 0: splat worker 8/16 reached 1024 local triangles.
Frame 0: splat worker 15/16 reached 2048 local triangles.
Frame 0: splat worker 9/16 reached 2048 local triangles.
Frame 0: splat worker 10/16 reached 2048 local triangles.
Frame 0: splat worker 2/16 reached 2048 local triangles.
Frame 0: splat worker 13/16 reached 1024 local triangles.
Frame 0: splat worker 1/16 reached 2048 local triangles.
Frame 0: splat worker 15/16 reached 4096 local triangles.
Frame 0: splat worker 5/16 reached 4096 local triangles.
Frame 0: splat worker 4/16 reached 4096 local triangles.
Frame 0: splat worker 6/16 reached 4096 local triangles.
Frame 0: splat worker 7/16 reached 2048 local triangles.
Frame 0: splat worker 11/16 reached 4096 local triangles.
Frame 0: splat worker 16/16 reached 4096 local triangles.
Frame 0: splat worker 3/16 reached 8192 local triangles.
Frame 0: splat worker 12/16 reached 4096 local triangles.
Frame 0: splat worker 14/16 reached 4096 local triangles.
Frame 0: splat worker 2/16 reached 4096 local triangles.
Frame 0: splat worker 13/16 reached 2048 local triangles.
Frame 0: splat worker 10/16 reached 4096 local triangles.
Frame 0: splat worker 8/16 reached 2048 local triangles.
Frame 0: splat worker 9/16 reached 4096 local triangles.
Frame 0: splat worker 1/16 reached 4096 local triangles.
Frame 0: splat worker 7/16 reached 4096 local triangles.
Frame 0: splat worker 15/16 reached 8192 local triangles.
Frame 0: splat worker 4/16 reached 8192 local triangles.
Frame 0: splat worker 10/16 reached 8192 local triangles.
Frame 0: splat worker 2/16 reached 8192 local triangles.
Frame 0: splat worker 8/16 reached 4096 local triangles.
Frame 0: splat worker 13/16 reached 4096 local triangles.
Frame 0: splat worker 11/16 reached 8192 local triangles.
Frame 0: splat worker 16/16 reached 8192 local triangles.
Frame 0: splat worker 1/16 finished after 7660 local triangles.
Frame 0: splat worker 10/16 finished after 9346 local triangles.
Frame 0: splat worker 4/16 finished after 9445 local triangles.
Frame 0: splat worker 13/16 finished after 4644 local triangles.
Frame 0: splat worker 11/16 finished after 8415 local triangles.
Frame 0: splat worker 12/16 finished after 7133 local triangles.
Frame 0: splat worker 14/16 finished after 7809 local triangles.
Frame 0: splat worker 16/16 finished after 8215 local triangles.
Frame 0: splat worker 15/16 finished after 10883 local triangles.
Frame 0: splat worker 9/16 finished after 6522 local triangles.
Frame 0: splat worker 5/16 finished after 8082 local triangles.
Frame 0: splat worker 7/16 finished after 6345 local triangles.
Frame 0: splat worker 2/16 finished after 8812 local triangles.
Frame 0: splat worker 3/16 finished after 14980 local triangles.
Frame 0: splat worker 6/16 finished after 7718 local triangles.
Frame 0: splat worker 8/16 finished after 4873 local triangles.
HIP producers: 16 workers; summed generation 2.053 s, submission/fallback 1.157 s; wall 0.326 s
Frame 0: catalog splatting finished in 0.4 s; writing image...
Rendered 156837 images from 21087 catalog stars to /tmp/gr-hip-final-multi.png (ok; imported lens map)
PSF splats: cached 156837, cached wing-clipped 0, direct fallbacks 0, discarded below min-Y 0
HIP PSF: 156837 events in 16 batches (16 timed); upload 0.000565 s, kernel 0.168641 s, download 0.027545 s
OMP_NUM_THREADS=16 timeout --kill-after=5s 60s --all-sky-catalog 4 1e-8 13.75s user 0.26s system 941% cpu 1.487 total
benchmarks/hip_psf_2026-09-07/final_multibatch.log
+3
View File
@@ -0,0 +1,3 @@
PSF cache ready: 64x64 phases, radius 47 px, relative tail 1e-08, tail abs 1e-06, boundary 1e-07, build 0.671 s
device=AMD Radeon RX 9070 arch=gfx1201 events=65536 spread=16 frame=3840x2160 event_bytes=56 CPU_reference=0.842402 s
HIP create=0.031723 replay_wall=0.083762 upload=0.000203 kernel=0.078268 download=0.005152 batches=4 timed=4 max_abs=1.7053025658242404e-13 max_rel=3.3472443502765663e-15 flux_rel=-1.3316774058238102e-17
benchmarks/hip_psf_2026-09-07/fixture_calibration_failed.log
+13
View File
@@ -0,0 +1,13 @@
RUN ['/home/wyj/Code/C/GR_4d_raytracing/build/Release/minkowski_sky', '--catalog', '/home/wyj/Code/C/GR_4d_raytracing/tests/data/hip_psf_mixed_catalog.csv', '--width', '96', '--height', '64', '--fov-deg', '30', '--look-ra-deg', '2', '--look-dec-deg', '1', '--coarse-cell-pixels', '16', '--refine-max-level', '0', '--exposure', '2', '--psf-fwhm-pixels', '2.7', '--psf-moffat-beta', '4.5', '--psf-relative-tail', '1e-8', '--max-cache-psf-flux', '1', '--psf-min-y', '0', '--hdr-output', '--output', '/tmp/gr-hip-renderer-test-z25me3rh/cpu-0.png']
PSF cache ready: 64x64 phases, radius 47 px, relative tail 1e-08, tail abs 1e-06, boundary 1e-07, build 1.641 s
Rendered 3 images from 3 catalog stars to /tmp/gr-hip-renderer-test-z25me3rh/cpu-0.png (ok)
PSF splats: cached 3, cached wing-clipped 0, direct fallbacks 0, discarded below min-Y 0
Traceback (most recent call last):
File "/home/wyj/Code/C/GR_4d_raytracing/tests/test_hip_renderer.py", line 61, in <module>
main()
~~~~^^
File "/home/wyj/Code/C/GR_4d_raytracing/tests/test_hip_renderer.py", line 47, in main
assert (stats[0], stats[2], stats[3]) == expected, stats
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
AssertionError: (3, 0, 0, 0)
benchmarks/hip_psf_2026-09-07/hdr_comparison.log
+9
View File
@@ -0,0 +1,9 @@
gr-hip-baseline-8192 gr-hip-final-8192 shape (3840, 2160, 3) samples 24883200 HDR_payload_exact True candidate_payload_sha256 423da7b025bf4815a5044e5d81c2f2330138cb8536b8058c19c9cba2d50f2f65
PNG_exact True
gr-hip-baseline-indexed gr-hip-final-indexed shape (3840, 2160, 3) samples 24883200 HDR_payload_exact True candidate_payload_sha256 423da7b025bf4815a5044e5d81c2f2330138cb8536b8058c19c9cba2d50f2f65
PNG_exact True
gr-hip-baseline-multi gr-hip-final-multi shape (3840, 2160, 3) samples 24883200 HDR_payload_exact True candidate_payload_sha256 a935becded15185af9e680aee9923ba51ecfee451a6bba34634a7e8be9a22ab8
PNG_exact True
gr-cpu-multi gr-hip-final-multi shape (3840, 2160, 3) samples 24883200 HDR_payload_exact True candidate_payload_sha256 a935becded15185af9e680aee9923ba51ecfee451a6bba34634a7e8be9a22ab8
PNG_exact True
clean baseline vs final: HDR_payload_exact True
benchmarks/hip_psf_2026-09-07/input_manifest.log
+3
View File
@@ -0,0 +1,3 @@
tile_ra252_dec048.csv source_stars 54160 sampled 8192 source_sha256 6ece2dc86b09709f3292a9d20f49d04cb53d520bd1b7d949f3a81306229ab48d sample_sha256 7c5516ff9ae92d0cfc8acfd67f578139a1ea2d3e2ca260cf9a382c38730d5095
tile_ra262_dec060.csv source_stars 76933 sampled 8192 source_sha256 9cc74db30205fe63fe8feb5581abc892efad44a079499f9185c2b162861b0bd7 sample_sha256 f6c9b4b0db950a660102497376778981322ce95187b1311ae515070fd6a5b88f
tile_ra082_dec090.csv source_stars 4703 sampled 4703 source_sha256 36f0713ac3c3d7cd849b5886a40228d1c535864e390cc5c055c80534aeef7b65 sample_sha256 36f0713ac3c3d7cd849b5886a40228d1c535864e390cc5c055c80534aeef7b65
benchmarks/hip_psf_2026-09-07/metadata.log
+31
View File
@@ -0,0 +1,31 @@
{
"base_commit": "265d7b95d5b916ec4768d174d73ad963a43f0088",
"optimized_source": "uncommitted worktree changes",
"compiler": "HIP version: 7.2.53210-9999\nclang version 22.1.8\nTarget: x86_64-pc-linux-gnu\nThread model: posix\nInstalledDir: /usr/lib/llvm/22/bin\nConfiguration file: /etc/clang/22/x86_64-pc-linux-gnu-clang++.cfg",
"hip_config_version": "7.2.53210-9999",
"kernel": "Linux 6.18.41-gentoo",
"cpu": [
"CPU(s): 16",
"Model name: 12th Gen Intel(R) Core(TM) i7-12700K",
"Thread(s) per core: 2",
"Core(s) per socket: 8",
"NUMA node0 CPU(s): 0-15"
],
"OMP_NUM_THREADS": 16,
"catalog_workers": 4,
"hip_flags": "--offload-arch=gfx1201 -std=c++17 -O2 -Wall -Wextra -Wpedantic",
"cpu_flags": "-std=c11 -march=native -pipe -Wall -Wextra -Wpedantic -O2 -DNDEBUG -fopenmp",
"event_bytes": 56,
"event_capacity": 16384,
"device_staging_slots": 2,
"sha256": {
"src/frame.c": "338d24eaa89d2d4c1e5a7f34d08ca827db13bb79ac3a6caeb155a731390ab023",
"src/hip_psf.hip": "b93fed04034a97e1eb25f3fc234da995287bf0c3fcde50609f95b0b3f4052a75",
"src/hip_psf.h": "00502f3d95c30d4c3d346e49dcf0e3cb846b5d461c2bfd5b3f24c7257791fad0",
"tests/benchmark_hip_psf.hip": "56cccf515b66762acb824556681c9f6986c0ffe3e3346eeadac23bc6d5926aa4",
"build/Release/schwarzschild_sky_hip": "1c340e1d8ab008e892101b46132069255ab92eb0d200bd91b2b300a4f4dacf6e",
"/tmp/gr-schwarzschild-hip-baseline": "de945e3dc826d559e1f8218a56d24f524c06e5382c8aa43667f9ac1fe5d2d719",
"/tmp/gr-hip-baseline-src/build/Release/schwarzschild_sky_hip": "de945e3dc826d559e1f8218a56d24f524c06e5382c8aa43667f9ac1fe5d2d719",
"output/lens/schwarzschild_galactic_center_R100_45deg_16_4_4k_.grlens": "905f960aea2d276f9586c79fba545a72cf742311e9e7e18faea293835c68bf28"
}
}
benchmarks/hip_psf_2026-09-07/prepare_inputs.py
+27
View File
@@ -0,0 +1,27 @@
#!/usr/bin/env python3
"""Recreate the bounded 2MASS samples used in this investigation."""
import hashlib
from pathlib import Path
repo = Path(__file__).resolve().parents[2]
source = repo / "assets/2mass/processed/all_sky"
multi = Path("/tmp/gr-hip-catalog-multi")
single = Path("/tmp/gr-hip-catalog-subset")
multi.mkdir(exist_ok=True)
single.mkdir(exist_ok=True)
for name in ["tile_ra252_dec048.csv", "tile_ra262_dec060.csv", "tile_ra082_dec090.csv"]:
raw = (source / name).read_bytes()
rows = raw.decode().splitlines()
n = min(8192, len(rows) - 1)
sample = ("\n".join([rows[0]] + [rows[1 + i * (len(rows) - 1) // n]
for i in range(n)]) + "\n").encode()
(multi / name).write_bytes(sample)
if name == "tile_ra252_dec048.csv":
(single / name).write_bytes(sample)
(Path(__file__).parent / "catalog_8192.csv").write_bytes(sample)
small = "\n".join([rows[0]] + [rows[1 + i * (len(rows) - 1) // 256]
for i in range(256)]) + "\n"
(Path(__file__).parent / "catalog_256.csv").write_text(small)
print(name, "source_stars", len(rows) - 1, "sampled", n,
"source_sha256", hashlib.sha256(raw).hexdigest(),
"sample_sha256", hashlib.sha256(sample).hexdigest())
benchmarks/hip_psf_2026-09-07/renderer_regression.log
+27
View File
@@ -0,0 +1,27 @@
RUN ['/home/wyj/Code/C/GR_4d_raytracing/build/Release/minkowski_sky', '--catalog', '/home/wyj/Code/C/GR_4d_raytracing/tests/data/hip_psf_mixed_catalog.csv', '--width', '96', '--height', '64', '--fov-deg', '30', '--look-ra-deg', '2', '--look-dec-deg', '1', '--coarse-cell-pixels', '16', '--refine-max-level', '0', '--exposure', '1000', '--psf-fwhm-pixels', '2.7', '--psf-moffat-beta', '4.5', '--psf-relative-tail', '1e-8', '--max-cache-psf-flux', '1', '--psf-min-y', '0', '--hdr-output', '--output', '/tmp/gr-hip-renderer-test-s9nk0w12/cpu-0.png']
PSF cache ready: 64x64 phases, radius 47 px, relative tail 1e-08, tail abs 1e-06, boundary 1e-07, build 1.667 s
Rendered 3 images from 3 catalog stars to /tmp/gr-hip-renderer-test-s9nk0w12/cpu-0.png (ok)
PSF splats: cached 2, cached wing-clipped 0, direct fallbacks 1, discarded below min-Y 0
RUN ['/home/wyj/Code/C/GR_4d_raytracing/build/Release/minkowski_sky_hip', '--catalog', '/home/wyj/Code/C/GR_4d_raytracing/tests/data/hip_psf_mixed_catalog.csv', '--width', '96', '--height', '64', '--fov-deg', '30', '--look-ra-deg', '2', '--look-dec-deg', '1', '--coarse-cell-pixels', '16', '--refine-max-level', '0', '--exposure', '1000', '--psf-fwhm-pixels', '2.7', '--psf-moffat-beta', '4.5', '--psf-relative-tail', '1e-8', '--max-cache-psf-flux', '1', '--psf-min-y', '0', '--hdr-output', '--output', '/tmp/gr-hip-renderer-test-s9nk0w12/hip-0.png']
PSF cache ready: 64x64 phases, radius 47 px, relative tail 1e-08, tail abs 1e-06, boundary 1e-07, build 1.700 s
HIP producers: 4 workers; summed generation 0.000 s, submission/fallback 0.020 s; wall 0.007 s
Rendered 3 images from 3 catalog stars to /tmp/gr-hip-renderer-test-s9nk0w12/hip-0.png (ok)
PSF splats: cached 2, cached wing-clipped 0, direct fallbacks 1, discarded below min-Y 0
HIP PSF: 2 events in 2 batches (2 timed); upload 0.000029 s, kernel 0.000887 s, download 0.000107 s
min_y=0: statistics match; HDR max_abs=0.0; exact=True
RUN ['/home/wyj/Code/C/GR_4d_raytracing/build/Release/minkowski_sky', '--catalog', '/home/wyj/Code/C/GR_4d_raytracing/tests/data/hip_psf_mixed_catalog.csv', '--width', '96', '--height', '64', '--fov-deg', '30', '--look-ra-deg', '2', '--look-dec-deg', '1', '--coarse-cell-pixels', '16', '--refine-max-level', '0', '--exposure', '1000', '--psf-fwhm-pixels', '2.7', '--psf-moffat-beta', '4.5', '--psf-relative-tail', '1e-8', '--max-cache-psf-flux', '1', '--psf-min-y', '1e30', '--hdr-output', '--output', '/tmp/gr-hip-renderer-test-s9nk0w12/cpu-1e30.png']
PSF cache ready: 64x64 phases, radius 47 px, relative tail 1e-08, tail abs 1e-06, boundary 1e-07, build 1.655 s
Rendered 3 images from 3 catalog stars to /tmp/gr-hip-renderer-test-s9nk0w12/cpu-1e30.png (ok)
PSF splats: cached 0, cached wing-clipped 0, direct fallbacks 0, discarded below min-Y 3
Warning: --psf-min-y discarded one or more PSF events.
RUN ['/home/wyj/Code/C/GR_4d_raytracing/build/Release/minkowski_sky_hip', '--catalog', '/home/wyj/Code/C/GR_4d_raytracing/tests/data/hip_psf_mixed_catalog.csv', '--width', '96', '--height', '64', '--fov-deg', '30', '--look-ra-deg', '2', '--look-dec-deg', '1', '--coarse-cell-pixels', '16', '--refine-max-level', '0', '--exposure', '1000', '--psf-fwhm-pixels', '2.7', '--psf-moffat-beta', '4.5', '--psf-relative-tail', '1e-8', '--max-cache-psf-flux', '1', '--psf-min-y', '1e30', '--hdr-output', '--output', '/tmp/gr-hip-renderer-test-s9nk0w12/hip-1e30.png']
PSF cache ready: 64x64 phases, radius 47 px, relative tail 1e-08, tail abs 1e-06, boundary 1e-07, build 1.664 s
HIP producers: 4 workers; summed generation 0.000 s, submission/fallback 0.000 s; wall 0.002 s
Rendered 3 images from 3 catalog stars to /tmp/gr-hip-renderer-test-s9nk0w12/hip-1e30.png (ok)
PSF splats: cached 0, cached wing-clipped 0, direct fallbacks 0, discarded below min-Y 3
Warning: --psf-min-y discarded one or more PSF events.
min_y=1e30: statistics match; HDR max_abs=0.0; exact=True
benchmarks/hip_psf_2026-09-07/streaming_test.log
+1
View File
@@ -0,0 +1 @@
HIP PSF cache comparison: max_abs=2.8421709430404007e-14 max_rel=7.5319184269262976e-16
benchmarks/hip_psf_2026-09-07/streaming_test_final.log
+1
View File
@@ -0,0 +1 @@
HIP PSF cache comparison: max_abs=2.8421709430404007e-14 max_rel=7.5319184269262976e-16
benchmarks/hip_psf_2026-09-07/wave_dense.log
+3
View File
@@ -0,0 +1,3 @@
PSF cache ready: 64x64 phases, radius 47 px, relative tail 1e-08, tail abs 1e-06, boundary 1e-07, build 0.691 s
device=AMD Radeon RX 9070 arch=gfx1201 events=4096 spread=16 frame=3840x2160 event_bytes=56 CPU_reference=0.051886 s
HIP create=0.023212 replay_wall=0.010834 upload=0.000021 kernel=0.005702 download=0.004809 batches=1 timed=1 max_abs=7.9936057773011271e-15 max_rel=2.9244848757288279e-15 flux_rel=3.55416413758235e-17
benchmarks/hip_psf_2026-09-07/wave_hotspot.log
+3
View File
@@ -0,0 +1,3 @@
PSF cache ready: 64x64 phases, radius 47 px, relative tail 1e-08, tail abs 1e-06, boundary 1e-07, build 0.671 s
device=AMD Radeon RX 9070 arch=gfx1201 events=16384 spread=1 frame=3840x2160 event_bytes=56 CPU_reference=0.208507 s
HIP create=0.024260 replay_wall=0.029202 upload=0.000050 kernel=0.023976 download=0.004869 batches=1 timed=1 max_abs=5.1159076974727213e-13 max_rel=4.4106699745012541e-15 flux_rel=3.1370299104282862e-18
benchmarks/hip_psf_2026-09-07/wave_spread.log
+3
View File
@@ -0,0 +1,3 @@
PSF cache ready: 64x64 phases, radius 47 px, relative tail 1e-08, tail abs 1e-06, boundary 1e-07, build 0.703 s
device=AMD Radeon RX 9070 arch=gfx1201 events=16384 spread=3840 frame=3840x2160 event_bytes=56 CPU_reference=0.422135 s
HIP create=0.025980 replay_wall=0.026740 upload=0.000047 kernel=0.021492 download=0.004899 batches=1 timed=1 max_abs=2.7755575615628914e-17 max_rel=5.6947355265748357e-16 flux_rel=4.0906184864912441e-20
build.md
+35
View File
@@ -89,6 +89,41 @@ mapping remain on the CPU. PSFs requiring direct evaluation retain the CPU
reference path with ordered HDR transfers. HIP initialization or execution
errors terminate the render instead of silently selecting the CPU backend.
HIP catalog mapping uses OpenMP producers (`OMP_NUM_THREADS`) with one private
16,384-event chunk per worker. A shared GPU sink owns one double HDR image and
PSF cache, plus two pinned upload slots protected by completion events. Each
PSF uses a group of 32 threads to visit adjacent pixels; weights, support and
HDR precision remain unchanged. CPU direct fallbacks exclude other submissions
through the complete download/evaluate/upload boundary.
Every batch contributes to the final HIP timings. Long submissions print
progress about every five seconds (completed counts cover reclaimed slots).
The producer summary reports **summed worker wall time** in generation and
submission/fallback, including lock/device waits in the latter; these overlapping
times must not be added to GPU timings or interpreted as process CPU time.
The two upload slots bound staging memory and allow CPU production ahead of GPU
completion; copies and kernels still execute in order on one stream.
With both Minkowski binaries built using `ENABLE_HDR=1`,
`python3 tests/test_hip_renderer.py build/Release/minkowski_sky build/Release/minkowski_sky_hip`
runs small CPU/HIP fallback and min-Y comparisons sequentially. Each renderer
subprocess has a 30-second timeout.
For a bounded production-cache benchmark, build and then run **one process at
a time**, with a timeout:
```sh
make PSF_BACKEND=hip SPACETIME=minkowski hip-psf-bench
./build/Release/benchmark_hip_psf --help
OMP_NUM_THREADS=16 timeout --kill-after=5s 45s ./build/Release/benchmark_hip_psf 4096 16
```
Arguments are event count (default 4096, maximum 65536) and screen spread in
pixels (default 16, maximum 3840). The benchmark runs a serial CPU HDR reference
before HIP and checks double values. See the [bounded investigation and original
logs](benchmarks/hip_psf_2026-09-07.md); its speedups do not establish a full-sky
render time.
## Rebuilding and checks
Build type, spacetime, HDR, and PSF backend configurations have separate object
gpu_psf_acceleration_plan.md
+15
View File
@@ -1,5 +1,20 @@
# GPU PSF 累积计划
## 2026-09-07 更新
有界实测后，production HIP 已改为 32-thread/event 协作 kernel，并恢复 OpenMP
catalog/inverse-map/颜色计算。每个 worker 拥有一个 16,384-event 私有块，共享的
GPU sink 只在批量提交和完整 CPU direct-fallback 边界加锁，保留单份 cache/HDR。
两个 pinned host/device slot 通过 kernel completion event 控制复用；所有 batch
均计时，长任务会周期报告已提交与已回收的完成事件数。详细结果、限制、重现命令及
原始输出见 [有界排查记录](benchmarks/hip_psf_2026-09-07.md)。
三种受控事件分布支持协作 kernel 的收益；真实三 tile 子集也改善了端到端耗时。
没有据此把全部 kernel 时间归因为 atomic contention，未引入 tile-local reduction、
float HDR 或黑体近似。下面第 7 节的 2026-09-05 状态是历史基线，单线程 producer
和前 64 批计时限制已被此次实现取代。实际 `PsfCachedEvent` 为 56 B，16,384 项为
896 KiB；早期按 24–32 B 估算的整帧内存不是当前结构的实际成本。
## 1. 目的与边界
本文件规划的是**最终点源 PSF 到 HDR framebuffer 的累积**加速，而不是 GPU
nr_spacetime_movie_renderer_design.md
+19
View File
@@ -958,6 +958,25 @@ while (!queue_empty()) {
---
## 22.3 可选 HIP PSF 后端
CPU reference 验证后，HIP 只替换最终 PSF 累积；catalog/source-triangle 查询、
inverse mapping、频移和黑体积分仍由 OpenMP worker 按 triangle 粗粒度并行完成。
每 worker 独占最多 16,384 项的 event chunk，不创建 per-star task。GPU sink 独占
一份 immutable cache 和 double RGB HDR；仅在 chunk 提交和 direct fallback 边界
使用共享锁。direct fallback 的完成 GPU 工作、下载、CPU 累积、上传必须在同一锁内。
每个 PSF 用 32 个线程协作访问相邻像素，仍使用原来的四点 phase cache 插值、圆形
support 和 double atomic 累积。两个 pinned host/device staging slot 只有在先前
kernel 完成后才能复用；所有 GPU 工作保持单 stream 顺序，不要求整个帧事件驻留内存。
计时 event 随 slot 复用，覆盖所有 batch，不再只测前 64 批。
以上参数由 [RX 9070 有界实验](benchmarks/hip_psf_2026-09-07.md) 支持，尚不是跨设备
最优值。该实验未运行完整全天渲染，不能用缩小样本直接宣称整帧加速比。分桶规约、
降低 HDR 精度或近似黑体积分均未作为此次优化的一部分。
---
# 23. Metric evaluator 的 thread-local workspace
nmesh 单点查询需要：
src/frame.c
+133 -22
View File
@@ -133,6 +133,9 @@ typedef struct {
const PsfKernelCache *cache;
#ifdef PSF_BACKEND_HIP
HipPsfSink *hip;
omp_lock_t *hip_lock; /* Shared sink lock; acquired only per chunk/fallback. */
int borrowed_hip;
double submit_seconds;
char hip_message[256];
HipPsfTiming hip_timing;
#endif
@@ -176,7 +179,7 @@ static int psf_event_sink_init(PsfEventSink *sink, double *hdr, int width,
return 0;
}
static int psf_event_sink_flush(PsfEventSink *sink) {
static int psf_event_sink_flush_unlocked(PsfEventSink *sink) {
if (sink->failed)
return -1;
#ifdef PSF_BACKEND_HIP
@@ -197,8 +200,27 @@ static int psf_event_sink_flush(PsfEventSink *sink) {
return 0;
}
static int psf_event_sink_flush(PsfEventSink *sink) {
#ifdef PSF_BACKEND_HIP
const double start = omp_get_wtime();
if (sink->hip_lock) omp_set_lock(sink->hip_lock);
#endif
const int result = psf_event_sink_flush_unlocked(sink);
#ifdef PSF_BACKEND_HIP
if (sink->hip_lock) omp_unset_lock(sink->hip_lock);
sink->submit_seconds += omp_get_wtime() - start;
#endif
return result;
}
/* A borrowed sink calls this only while holding hip_lock across the entire
* download -> CPU fallback -> upload transaction. */
static int psf_event_sink_finish_for_cpu(PsfEventSink *sink) {
#ifdef PSF_BACKEND_HIP
if (psf_event_sink_flush_unlocked(sink))
#else
if (psf_event_sink_flush(sink))
#endif
return -1;
#ifdef PSF_BACKEND_HIP
if (sink->hip != NULL && hip_psf_sink_finish(sink->hip, sink->hdr,
@@ -226,6 +248,13 @@ static int psf_event_sink_resume_gpu(PsfEventSink *sink) {
}
static int psf_event_sink_destroy(PsfEventSink *sink) {
#ifdef PSF_BACKEND_HIP
if (sink->borrowed_hip) {
const int result = psf_event_sink_flush(sink);
free(sink->events);
return result;
}
#endif
int result = psf_event_sink_finish_for_cpu(sink);
#ifdef PSF_BACKEND_HIP
if (sink->hip != NULL && hip_psf_sink_get_timing(sink->hip, &sink->hip_timing)) {
@@ -238,7 +267,7 @@ static int psf_event_sink_destroy(PsfEventSink *sink) {
return result;
}
#if FRAME_PSF_EVENT_SINK
#if FRAME_PSF_EVENT_SINK || defined(PSF_BACKEND_HIP)
static void psf_event_sink_emit(PsfEventSink *sink, const PsfCachedEvent *event) {
if (sink->failed)
return;
@@ -1086,16 +1115,27 @@ static int splat_catalog_tile(const Star *stars, size_t count,
&event, image_x, image_y, color, flux, context->psf, context->psf_cache,
context->max_cache_psf_flux, context->psf_relative_tail, context->psf_min_y);
if (direct_fallback == 1) {
/* A direct fallback forms an explicit ordered CPU/GPU boundary. */
if (psf_event_sink_finish_for_cpu(context->event_sink))
return -1;
splat_moffat_direct(context->hdr, context->width, context->height, image_x,
image_y, color, flux, context->psf,
context->psf_relative_tail, context->psf_min_y);
if (psf_event_sink_resume_gpu(context->event_sink))
return -1;
/* Exclude every other submit and fallback until HDR is back on device. */
#ifdef PSF_BACKEND_HIP
const double submit_start = omp_get_wtime();
if (context->event_sink->hip_lock)
omp_set_lock(context->event_sink->hip_lock);
#endif
int failed = psf_event_sink_finish_for_cpu(context->event_sink);
if (!failed) {
splat_moffat_direct(context->hdr, context->width, context->height, image_x,
image_y, color, flux, context->psf,
context->psf_relative_tail, context->psf_min_y);
failed = psf_event_sink_resume_gpu(context->event_sink);
}
#ifdef PSF_BACKEND_HIP
if (context->event_sink->hip_lock)
omp_unset_lock(context->event_sink->hip_lock);
context->event_sink->submit_seconds += omp_get_wtime() - submit_start;
#endif
if (failed) return -1;
} else if (direct_fallback != 3) {
#if FRAME_PSF_EVENT_SINK
#if FRAME_PSF_EVENT_SINK || defined(PSF_BACKEND_HIP)
psf_event_sink_emit(context->event_sink, &event);
#else
splat_prepared_cached_event(context->hdr, context->width, context->height,
@@ -1261,20 +1301,91 @@ size_t frame_splat_catalog(const FrameLensMesh *mesh,
mesh->triangle_count);
#ifdef PSF_BACKEND_HIP
/* A single GPU HDR framebuffer owns all cached events. Keep catalog/lens
* work serial for this first direct-atomic integration; the CPU parallel
* private-HDR path remains the PSF_BACKEND=cpu implementation. */
const CatalogSplatStats hip_stats = splat_catalog_triangles(
mesh, catalog, hdr, width, height, exposure, psf, psf_cache,
max_magnification, max_cache_psf_flux, psf_relative_tail, psf_min_y,
0, mesh->triangle_count, NULL);
copy_psf_splat_stats(psf_stats, hip_stats);
if (hip_stats.failed)
return SIZE_MAX;
/* CPU workers own only bounded event chunks. A single device cache/HDR is
* borrowed under a coarse submission lock, never duplicated per worker. */
PsfEventSink owner;
if (psf_event_sink_init(&owner, hdr, width, height, psf_cache)) return SIZE_MAX;
omp_lock_t submit_lock;
omp_init_lock(&submit_lock);
size_t hip_images = 0, hip_direct = 0, hip_clipped = 0, hip_discarded = 0;
int hip_failed = 0, hip_workers = 0;
double hip_generate_seconds = 0.0, hip_submit_seconds = 0.0;
const double hip_start = omp_get_wtime();
#ifdef GR_DEBUG
double hip_max_magnification = 0.0;
size_t hip_clamped = 0;
#pragma omp parallel reduction(+ : hip_images, hip_direct, hip_clipped, hip_discarded, hip_clamped, hip_generate_seconds, hip_submit_seconds) reduction(max : hip_failed, hip_max_magnification)
#else
#pragma omp parallel reduction(+ : hip_images, hip_direct, hip_clipped, hip_discarded, hip_generate_seconds, hip_submit_seconds) reduction(max : hip_failed)
#endif
{
const double worker_start = omp_get_wtime();
#pragma omp single
hip_workers = omp_get_num_threads();
PsfEventSink local = {.hdr = hdr, .width = width, .height = height,
.cache = psf_cache, .hip = owner.hip, .hip_lock = &submit_lock,
.borrowed_hip = 1};
local.events = malloc(PSF_EVENT_SINK_CAPACITY * sizeof *local.events);
if (!local.events) {
snprintf(local.hip_message, sizeof local.hip_message, "worker event allocation failed");
psf_event_sink_mark_failed(&local, "producer initialization");
}
const size_t worker = (size_t)omp_get_thread_num();
const size_t workers = (size_t)omp_get_num_threads();
size_t completed = 0, next_report = 8;
if (progress && progress->worker_callback)
progress->worker_callback(progress->context, worker, workers, 0, 0);
#pragma omp for schedule(dynamic, 1) nowait
for (size_t t = 0; t < mesh->triangle_count; ++t) {
if (local.failed) continue;
const CatalogSplatStats stats = splat_catalog_triangles(
mesh, catalog, hdr, width, height, exposure, psf, psf_cache,
max_magnification, max_cache_psf_flux, psf_relative_tail, psf_min_y,
t, t + 1, &local);
hip_images += stats.images;
hip_direct += stats.direct_fallbacks;
hip_clipped += stats.cached_wing_clipped;
hip_discarded += stats.discarded_below_min_y;
hip_failed |= stats.failed;
#ifdef GR_DEBUG
hip_max_magnification = fmax(hip_max_magnification, stats.max_raw_magnification);
hip_clamped += stats.magnification_clamped_triangles;
#endif
++completed;
if (progress && progress->worker_callback && completed == next_report) {
progress->worker_callback(progress->context, worker, workers, completed, 0);
if (next_report <= SIZE_MAX / 2) next_report *= 2;
}
}
if (psf_event_sink_destroy(&local)) hip_failed = 1;
hip_submit_seconds += local.submit_seconds;
hip_generate_seconds += omp_get_wtime() - worker_start - local.submit_seconds;
if (progress && progress->worker_callback)
progress->worker_callback(progress->context, worker, workers, completed, 1);
}
omp_destroy_lock(&submit_lock);
if (psf_event_sink_destroy(&owner)) hip_failed = 1;
fprintf(stderr, "HIP producers: %d workers; summed generation %.3f s, submission/fallback %.3f s; wall %.3f s\n",
hip_workers, hip_generate_seconds, hip_submit_seconds, omp_get_wtime() - hip_start);
copy_psf_splat_stats(psf_stats, (CatalogSplatStats){
.images = hip_images, .direct_fallbacks = hip_direct,
.cached_wing_clipped = hip_clipped, .discarded_below_min_y = hip_discarded,
.gpu_event_count = owner.hip_timing.event_count,
.gpu_batch_count = owner.hip_timing.batch_count,
.gpu_timed_batch_count = owner.hip_timing.timed_batch_count,
.gpu_upload_seconds = owner.hip_timing.upload_seconds,
.gpu_kernel_seconds = owner.hip_timing.kernel_seconds,
.gpu_download_seconds = owner.hip_timing.download_seconds,
#ifdef GR_DEBUG
.max_raw_magnification = hip_max_magnification,
.magnification_clamped_triangles = hip_clamped,
#endif
});
if (hip_failed) return SIZE_MAX;
if (progress != NULL && progress->callback != NULL)
progress->callback(progress->context, FRAME_SPLAT_PROGRESS_END,
mesh->triangle_count, mesh->triangle_count);
return hip_stats.images;
return hip_images;
#endif
const size_t pixel_count = (size_t)width * height * 3;
src/hip_psf.h
+6 -4
View File
@@ -32,16 +32,18 @@ int hip_psf_sink_create(HipPsfSink **sink, int width, int height,
int hip_psf_sink_load_hdr(HipPsfSink *sink, const double *hdr,
char *message, size_t message_size);
/* Queues a cached-event chunk. Direct fallbacks and min-Y discards stay with
* the caller; event_count must not exceed the creation capacity. */
/* Copies a chunk into one of two owned pinned staging slots; caller memory
* may be reused on return. Slot reuse waits for its previous kernel. Calls
* sharing a sink must be serialized by the owner. Direct fallbacks and min-Y
* discards stay with the caller; event_count must not exceed capacity. */
int hip_psf_sink_submit(HipPsfSink *sink, const PsfCachedEvent *events,
size_t event_count, char *message, size_t message_size);
/* Completes all work and overwrites hdr with double linear RGB device HDR. */
int hip_psf_sink_finish(HipPsfSink *sink, double *hdr, char *message,
size_t message_size);
/* Valid after hip_psf_sink_finish(); timings are accumulated without a
* per-batch host synchronization. */
/* Complete after finish(); otherwise includes recycled/completed slots only.
* Every batch is timed with a fixed-size reusable event pool. */
int hip_psf_sink_get_timing(const HipPsfSink *sink, HipPsfTiming *timing);
void hip_psf_sink_destroy(HipPsfSink *sink);
src/hip_psf.hip
+103 -64
View File
@@ -5,27 +5,34 @@
#include <cmath>
#include <cstdio>
#include <limits>
#include <vector>
#include <chrono>
#include <cstring>
struct HipPsfBatchTiming {
enum { PSF_LANES_PER_EVENT = 32, PSF_THREADS_PER_BLOCK = 128 };
/* Two reusable staging slots bound host/device queue memory. A slot may be
* overwritten only after its kernel_end event completes. All API calls for a
* sink are serialized by its owner, including calls from CPU producer workers. */
struct HipPsfBatch {
PsfCachedEvent *host = nullptr, *device = nullptr;
hipEvent_t upload_start = nullptr, upload_end = nullptr;
hipEvent_t kernel_start = nullptr, kernel_end = nullptr;
size_t count = 0;
};
enum { HIP_PSF_MAX_TIMED_BATCHES = 64 };
struct HipPsfSink {
PsfCachedEvent *events = nullptr;
HipPsfBatch slots[2];
float *weights = nullptr;
double *hdr = nullptr;
hipStream_t stream = nullptr;
size_t event_capacity = 0, hdr_values = 0;
size_t completed_events = 0;
int device = 0;
int width = 0, height = 0, phase_resolution = 0, radius_pixels = 0;
double max_radius_pixels = 0.0;
hipEvent_t download_start = nullptr, download_end = nullptr;
std::vector<HipPsfBatchTiming> batches;
size_t measured_batches = 0;
HipPsfTiming timing = {};
std::chrono::steady_clock::time_point last_report = std::chrono::steady_clock::now();
};
static int report(hipError_t status, char *message, size_t message_size) {
@@ -37,13 +44,19 @@ static void ok(char *message, size_t message_size) {
if (message && message_size) std::snprintf(message, message_size, "ok");
}
static void destroy_batch_timing(HipPsfBatchTiming *timing) {
if (!timing) return;
(void)hipEventDestroy(timing->upload_start);
(void)hipEventDestroy(timing->upload_end);
(void)hipEventDestroy(timing->kernel_start);
(void)hipEventDestroy(timing->kernel_end);
*timing = {};
static int collect_batch(HipPsfSink *sink, HipPsfBatch *batch,
char *message, size_t message_size) {
if (!batch->count) return 0;
float upload_ms = 0.0f, kernel_ms = 0.0f;
if (report(hipEventSynchronize(batch->kernel_end), message, message_size) ||
report(hipEventElapsedTime(&upload_ms, batch->upload_start, batch->upload_end), message, message_size) ||
report(hipEventElapsedTime(&kernel_ms, batch->kernel_start, batch->kernel_end), message, message_size)) return -1;
sink->timing.upload_seconds += upload_ms * 1e-3;
sink->timing.kernel_seconds += kernel_ms * 1e-3;
++sink->timing.timed_batch_count;
sink->completed_events += batch->count;
batch->count = 0;
return 0;
}
__device__ static size_t weight_index(int phase_resolution, int radius_pixels,
@@ -71,7 +84,11 @@ __global__ static void splat_kernel(const PsfCachedEvent *events, size_t count,
int width, int height, const float *weights,
int phase_resolution, int radius_pixels,
double max_radius_pixels, double *hdr) {
const size_t index = (size_t)blockIdx.x * blockDim.x + threadIdx.x;
/* Consecutive lanes visit consecutive pixels of the same phase plane.
* Fixed 32-lane groups also work on wave64 devices without wave intrinsics. */
const size_t thread = (size_t)blockIdx.x * blockDim.x + threadIdx.x;
const size_t index = thread / PSF_LANES_PER_EVENT;
const int lane = thread % PSF_LANES_PER_EVENT;
if (index >= count) return;
const PsfCachedEvent event = events[index];
const int base_x = (int)floor(event.x), base_y = (int)floor(event.y);
@@ -88,7 +105,7 @@ __global__ static void splat_kernel(const PsfCachedEvent *events, size_t count,
if (!row_range(event.support_radius, fx, fy, support, offset_y, &first, &last)) continue;
if (first < -base_x) first = -base_x;
if (last >= width - base_x) last = width - base_x - 1;
for (int offset_x = first; offset_x <= last; ++offset_x) {
for (int offset_x = first + lane; offset_x <= last; offset_x += PSF_LANES_PER_EVENT) {
const int px = base_x + offset_x;
const double w00 = weights[weight_index(phase_resolution, radius_pixels, x0, y0, offset_x, offset_y)];
const double w10 = weights[weight_index(phase_resolution, radius_pixels, x1, y0, offset_x, offset_y)];
@@ -119,15 +136,18 @@ extern "C" int hip_psf_sink_create(HipPsfSink **out, int width, int height,
const PsfKernelCache *cache, size_t event_capacity,
char *message, size_t message_size) {
if (!out || width <= 0 || height <= 0 || !cache || !cache->ready || !cache->weights ||
cache->phase_resolution <= 0 || cache->radius_pixels < 0 || event_capacity == 0) {
cache->phase_resolution <= 0 || cache->radius_pixels < 0 || event_capacity == 0 ||
event_capacity > std::numeric_limits<size_t>::max() / sizeof(PsfCachedEvent)) {
if (message && message_size) std::snprintf(message, message_size, "invalid HIP PSF sink arguments");
return -1;
}
const size_t nodes = (size_t)cache->phase_resolution + 1;
const size_t side = (size_t)cache->radius_pixels * 2 + 1;
if (nodes > std::numeric_limits<size_t>::max() / nodes || nodes * nodes > std::numeric_limits<size_t>::max() / side ||
nodes * nodes * side > std::numeric_limits<size_t>::max() / side || (size_t)width > std::numeric_limits<size_t>::max() / (size_t)height ||
(size_t)width * (size_t)height > std::numeric_limits<size_t>::max() / 3) {
nodes * nodes * side > std::numeric_limits<size_t>::max() / side ||
nodes * nodes * side * side > std::numeric_limits<size_t>::max() / sizeof(float) ||
(size_t)width > std::numeric_limits<size_t>::max() / (size_t)height ||
(size_t)width * (size_t)height > std::numeric_limits<size_t>::max() / (3 * sizeof(double))) {
if (message && message_size) std::snprintf(message, message_size, "HIP PSF sink size overflow");
return -1;
}
@@ -138,16 +158,30 @@ extern "C" int hip_psf_sink_create(HipPsfSink **out, int width, int height,
sink->width = width; sink->height = height; sink->phase_resolution = cache->phase_resolution;
sink->radius_pixels = cache->radius_pixels; sink->max_radius_pixels = cache->max_radius_pixels;
const size_t weight_count = nodes * nodes * side * side;
if (report(hipStreamCreate(&sink->stream), message, message_size) ||
if (report(hipGetDevice(&sink->device), message, message_size) ||
report(hipStreamCreate(&sink->stream), message, message_size) ||
report(hipEventCreate(&sink->download_start), message, message_size) ||
report(hipEventCreate(&sink->download_end), message, message_size) ||
report(hipMalloc(&sink->events, event_capacity * sizeof *sink->events), message, message_size) ||
report(hipMalloc(&sink->weights, weight_count * sizeof *sink->weights), message, message_size) ||
report(hipMalloc(&sink->hdr, sink->hdr_values * sizeof *sink->hdr), message, message_size) ||
report(hipMemcpyAsync(sink->weights, cache->weights, weight_count * sizeof *sink->weights, hipMemcpyHostToDevice, sink->stream), message, message_size) ||
report(hipMemsetAsync(sink->hdr, 0, sink->hdr_values * sizeof *sink->hdr, sink->stream), message, message_size)) {
hip_psf_sink_destroy(sink); return -1;
}
for (HipPsfBatch &batch : sink->slots) {
if (report(hipHostMalloc(&batch.host, event_capacity * sizeof(PsfCachedEvent)), message, message_size) ||
report(hipMalloc(&batch.device, event_capacity * sizeof(PsfCachedEvent)), message, message_size) ||
report(hipEventCreate(&batch.upload_start), message, message_size) ||
report(hipEventCreate(&batch.upload_end), message, message_size) ||
report(hipEventCreate(&batch.kernel_start), message, message_size) ||
report(hipEventCreate(&batch.kernel_end), message, message_size)) {
hip_psf_sink_destroy(sink); return -1;
}
}
/* The caller may release the immutable host cache after creation. */
if (report(hipStreamSynchronize(sink->stream), message, message_size)) {
hip_psf_sink_destroy(sink); return -1;
}
*out = sink; ok(message, message_size); return 0;
}
@@ -158,37 +192,37 @@ extern "C" int hip_psf_sink_submit(HipPsfSink *sink, const PsfCachedEvent *event
return -1;
}
if (!event_count) { ok(message, message_size); return 0; }
const unsigned int threads = 128;
const size_t blocks = (event_count + threads - 1) / threads;
HipPsfBatchTiming timing;
const int timed = sink->batches.size() < HIP_PSF_MAX_TIMED_BATCHES;
if (blocks > std::numeric_limits<unsigned int>::max() ||
(timed && (report(hipEventCreate(&timing.upload_start), message, message_size) ||
report(hipEventCreate(&timing.upload_end), message, message_size) ||
report(hipEventCreate(&timing.kernel_start), message, message_size) ||
report(hipEventCreate(&timing.kernel_end), message, message_size) ||
report(hipEventRecord(timing.upload_start, sink->stream), message, message_size))) ||
report(hipMemcpyAsync(sink->events, events, event_count * sizeof *events,
hipMemcpyHostToDevice, sink->stream), message, message_size) ||
(timed && (report(hipEventRecord(timing.upload_end, sink->stream), message, message_size) ||
report(hipEventRecord(timing.kernel_start, sink->stream), message, message_size)))) {
destroy_batch_timing(&timing);
if (report(hipSetDevice(sink->device), message, message_size)) return -1;
const unsigned int threads = PSF_THREADS_PER_BLOCK;
const size_t per_block = PSF_THREADS_PER_BLOCK / PSF_LANES_PER_EVENT;
const size_t blocks = event_count / per_block + (event_count % per_block != 0);
if (blocks > std::numeric_limits<unsigned int>::max()) {
if (message && message_size) std::snprintf(message, message_size, "HIP PSF grid too large");
return -1;
}
HipPsfBatch &batch = sink->slots[sink->timing.batch_count % 2];
if (collect_batch(sink, &batch, message, message_size)) return -1;
std::memcpy(batch.host, events, event_count * sizeof *events);
if (report(hipEventRecord(batch.upload_start, sink->stream), message, message_size) ||
report(hipMemcpyAsync(batch.device, batch.host, event_count * sizeof *events,
hipMemcpyHostToDevice, sink->stream), message, message_size) ||
report(hipEventRecord(batch.upload_end, sink->stream), message, message_size) ||
report(hipEventRecord(batch.kernel_start, sink->stream), message, message_size)) return -1;
hipLaunchKernelGGL(splat_kernel, dim3((unsigned int)blocks), dim3(threads), 0, sink->stream,
sink->events, event_count, sink->width, sink->height, sink->weights,
batch.device, event_count, sink->width, sink->height, sink->weights,
sink->phase_resolution, sink->radius_pixels, sink->max_radius_pixels, sink->hdr);
if (report(hipGetLastError(), message, message_size) ||
(timed && report(hipEventRecord(timing.kernel_end, sink->stream), message, message_size))) {
destroy_batch_timing(&timing);
return -1;
}
if (timed) {
sink->batches.push_back(timing);
++sink->timing.timed_batch_count;
}
report(hipEventRecord(batch.kernel_end, sink->stream), message, message_size)) return -1;
batch.count = event_count;
++sink->timing.batch_count;
sink->timing.event_count += event_count;
const auto now = std::chrono::steady_clock::now();
if (std::chrono::duration<double>(now - sink->last_report).count() >= 5.0) {
std::fprintf(stderr, "HIP PSF progress: %zu submitted / %zu completed events; %zu / %zu batches timed; kernel %.3f s\n",
sink->timing.event_count, sink->completed_events, sink->timing.timed_batch_count,
sink->timing.batch_count, sink->timing.kernel_seconds);
sink->last_report = now;
}
ok(message, message_size); return 0;
}
@@ -198,29 +232,24 @@ extern "C" int hip_psf_sink_load_hdr(HipPsfSink *sink, const double *hdr,
if (message && message_size) std::snprintf(message, message_size, "invalid HIP PSF HDR upload");
return -1;
}
if (report(hipMemcpyAsync(sink->hdr, hdr, sink->hdr_values * sizeof *hdr,
hipMemcpyHostToDevice, sink->stream), message, message_size))
if (report(hipSetDevice(sink->device), message, message_size) ||
report(hipMemcpyAsync(sink->hdr, hdr, sink->hdr_values * sizeof *hdr,
hipMemcpyHostToDevice, sink->stream), message, message_size) ||
report(hipStreamSynchronize(sink->stream), message, message_size))
return -1;
ok(message, message_size); return 0;
}
extern "C" int hip_psf_sink_finish(HipPsfSink *sink, double *hdr, char *message, size_t message_size) {
if (!sink || !hdr) { if (message && message_size) std::snprintf(message, message_size, "invalid HIP PSF HDR download"); return -1; }
if (report(hipEventRecord(sink->download_start, sink->stream), message, message_size) ||
if (report(hipSetDevice(sink->device), message, message_size) ||
report(hipEventRecord(sink->download_start, sink->stream), message, message_size) ||
report(hipMemcpyAsync(hdr, sink->hdr, sink->hdr_values * sizeof *hdr, hipMemcpyDeviceToHost, sink->stream), message, message_size) ||
report(hipEventRecord(sink->download_end, sink->stream), message, message_size) ||
report(hipStreamSynchronize(sink->stream), message, message_size)) return -1;
float milliseconds = 0.0f;
for (size_t i = sink->measured_batches; i < sink->batches.size(); ++i) {
const HipPsfBatchTiming &batch = sink->batches[i];
if (report(hipEventElapsedTime(&milliseconds, batch.upload_start, batch.upload_end),
message, message_size)) return -1;
sink->timing.upload_seconds += milliseconds * 1e-3;
if (report(hipEventElapsedTime(&milliseconds, batch.kernel_start, batch.kernel_end),
message, message_size)) return -1;
sink->timing.kernel_seconds += milliseconds * 1e-3;
}
sink->measured_batches = sink->batches.size();
for (HipPsfBatch &batch : sink->slots)
if (collect_batch(sink, &batch, message, message_size)) return -1;
if (report(hipEventElapsedTime(&milliseconds, sink->download_start, sink->download_end),
message, message_size)) return -1;
sink->timing.download_seconds += milliseconds * 1e-3;
@@ -235,10 +264,20 @@ extern "C" int hip_psf_sink_get_timing(const HipPsfSink *sink, HipPsfTiming *tim
extern "C" void hip_psf_sink_destroy(HipPsfSink *sink) {
if (!sink) return;
for (HipPsfBatchTiming &timing : sink->batches)
destroy_batch_timing(&timing);
(void)hipEventDestroy(sink->download_start);
(void)hipEventDestroy(sink->download_end);
(void)hipFree(sink->events); (void)hipFree(sink->weights); (void)hipFree(sink->hdr);
(void)hipStreamDestroy(sink->stream); delete sink;
(void)hipSetDevice(sink->device);
if (sink->stream) (void)hipStreamSynchronize(sink->stream);
for (HipPsfBatch &batch : sink->slots) {
if (batch.upload_start) (void)hipEventDestroy(batch.upload_start);
if (batch.upload_end) (void)hipEventDestroy(batch.upload_end);
if (batch.kernel_start) (void)hipEventDestroy(batch.kernel_start);
if (batch.kernel_end) (void)hipEventDestroy(batch.kernel_end);
if (batch.host) (void)hipHostFree(batch.host);
if (batch.device) (void)hipFree(batch.device);
}
if (sink->download_start) (void)hipEventDestroy(sink->download_start);
if (sink->download_end) (void)hipEventDestroy(sink->download_end);
if (sink->weights) (void)hipFree(sink->weights);
if (sink->hdr) (void)hipFree(sink->hdr);
if (sink->stream) (void)hipStreamDestroy(sink->stream);
delete sink;
}
tests/benchmark_hip_psf.hip
+96
View File
@@ -0,0 +1,96 @@
#include "hip_psf.h"
#include <hip/hip_runtime.h>
#include <omp.h>
#include <cmath>
#include <cstdio>
#include <cstdlib>
#include <cstring>
#include <vector>
/* Bounded production-cache replay. One process runs CPU reference then HIP;
* no catalog loading or geodesic work. Arguments: events, spread pixels. */
int main(int argc, char **argv) {
if (argc == 2 && !std::strcmp(argv[1], "--help")) {
std::puts("Usage: benchmark_hip_psf [EVENTS [SPREAD_PIXELS]]\n"
" EVENTS Cached events, 1..65536 (default: 4096)\n"
" SPREAD_PIXELS Star-center spread, 1..3840 (default: 16)\n"
"Runs a serial CPU reference, then HIP, at 3840x2160; PSF 2.7/4.5, tail 1e-8.\n"
"Use an external timeout and run only one benchmark process at a time.");
return 0;
}
auto number = [](const char *s, long upper) {
char *end = nullptr;
const long n = std::strtol(s, &end, 10);
return s != end && !*end && n >= 1 && n <= upper ? (int)n : 0;
};
const int count = argc > 1 ? number(argv[1], 65536) : 4096;
const int spread = argc > 2 ? number(argv[2], 3840) : 16;
if (argc > 3 || !count || !spread) {
std::fputs("Invalid bounded benchmark arguments; see --help.\n", stderr);
return 2;
}
std::setvbuf(stdout, nullptr, _IOLBF, 0);
const int width = 3840, height = 2160;
const size_t values = (size_t)width * height * 3;
PsfKernelCache cache = {};
const PointSpreadFunction psf = {2.7, 4.5};
if (psf_kernel_cache_init(&cache, &psf, 1e-8)) return 1;
psf_kernel_cache_report_ready(&cache, stdout);
std::vector<PsfCachedEvent> events(count);
unsigned state = 12345;
auto uniform = [&]() { state = state * 1664525u + 1013904223u;
return (state >> 8) * (1.0 / 16777216.0); };
for (auto &event : events) {
const double x = (width - spread) * 0.5 + uniform() * spread;
const double sy = spread < height ? spread : height;
const double y = (height - sy) * 0.5 + uniform() * sy;
const LinearRgb color = {0.1 + uniform(), 0.1 + uniform(), 0.1 + uniform()};
const int result = psf_prepare_cached_event(&event, x, y, color,
0.1 + uniform(), &psf, &cache, 1e8, 1e-8, 0);
if (result != 0) return 1;
}
std::vector<double> cpu(values, 0), gpu(values, 0);
double start = omp_get_wtime();
for (const auto &event : events)
splat_prepared_cached_event(cpu.data(), width, height, &event, &cache);
const double cpu_seconds = omp_get_wtime() - start;
char message[256] = {};
hipDeviceProp_t prop;
if (hipGetDeviceProperties(&prop, 0) != hipSuccess) return 1;
std::printf("device=%s arch=%s events=%d spread=%d frame=%dx%d event_bytes=%zu CPU_reference=%.6f s\n",
prop.name, prop.gcnArchName, count, spread, width, height, sizeof(PsfCachedEvent), cpu_seconds);
HipPsfSink *sink = nullptr;
start = omp_get_wtime();
if (hip_psf_sink_create(&sink, width, height, &cache, 16384, message, sizeof message)) {
std::fprintf(stderr, "create: %s\n", message); return 1;
}
const double create_seconds = omp_get_wtime() - start;
start = omp_get_wtime();
for (size_t i = 0; i < events.size(); i += 16384) {
size_t n = events.size() - i;
if (n > 16384) n = 16384;
if (hip_psf_sink_submit(sink, events.data() + i, n, message, sizeof message)) {
std::fprintf(stderr, "submit: %s\n", message); return 1;
}
}
if (hip_psf_sink_finish(sink, gpu.data(), message, sizeof message)) {
std::fprintf(stderr, "finish: %s\n", message); return 1;
}
const double wall = omp_get_wtime() - start;
HipPsfTiming timing = {};
hip_psf_sink_get_timing(sink, &timing);
double max_abs = 0, max_rel = 0, sum = 0, diff = 0;
for (size_t i = 0; i < values; ++i) {
if (!std::isfinite(gpu[i])) return 1;
const double d = std::fabs(gpu[i] - cpu[i]);
max_abs = std::fmax(max_abs, d);
if (cpu[i] != 0) max_rel = std::fmax(max_rel, d / std::fabs(cpu[i]));
sum += cpu[i]; diff += gpu[i] - cpu[i];
}
std::printf("HIP create=%.6f replay_wall=%.6f upload=%.6f kernel=%.6f download=%.6f batches=%zu timed=%zu max_abs=%.17g max_rel=%.17g flux_rel=%.17g\n",
create_seconds, wall, timing.upload_seconds, timing.kernel_seconds, timing.download_seconds,
timing.batch_count, timing.timed_batch_count, max_abs, max_rel, diff / sum);
hip_psf_sink_destroy(sink);
psf_kernel_cache_destroy(&cache);
return max_abs < 1e-10 && max_rel < 1e-10 ? 0 : 1;
}
tests/test_hip_psf.hip
+32 -1
View File
@@ -57,7 +57,38 @@ int main() {
return 1;
}
double max_abs = 0.0, max_rel = 0.0;
/* Exercise both reusable slots, more than the old 64-batch timing limit,
* immediate caller-buffer reuse, edge clipping and an ordered CPU boundary. */
for (int batch = 0; batch < 70; ++batch) {
for (int j = 0; j < 3; ++j) {
events[j] = PsfCachedEvent{};
const int result = psf_prepare_cached_event(&events[j],
-4.75 + (batch * 7 + j * 13) % 75,
-3.125 + (batch * 11 + j * 5) % 58, colors[j],
batch % 7 == 0 ? 1000.0 : 0.01, &psf, &cache, 1e8, 1e-8, 0.0);
if (result != 0 && result != 2) return 1;
splat_prepared_cached_event(cpu_hdr, width, height, &events[j], &cache);
}
if (hip_psf_sink_submit(sink, events, 3, message, sizeof message)) {
std::fprintf(stderr, "streaming submit failed: %s\n", message); return 1;
}
for (auto &event : events) event = PsfCachedEvent{};
if (batch == 35) {
if (hip_psf_sink_finish(sink, gpu_hdr, message, sizeof message)) return 1;
for (double *hdr : {cpu_hdr, gpu_hdr})
splat_moffat_direct(hdr, width, height, 12.25, 20.75, colors[0],
0.2, &psf, 1e-8, 0.0);
if (hip_psf_sink_load_hdr(sink, gpu_hdr, message, sizeof message)) return 1;
}
}
if (hip_psf_sink_finish(sink, gpu_hdr, message, sizeof message) ||
hip_psf_sink_get_timing(sink, &timing) || timing.event_count != 213 ||
timing.batch_count != 71 || timing.timed_batch_count != 71) {
std::fprintf(stderr, "streaming completion/accounting failed: %s\n", message);
return 1;
}
for (size_t i = 0; i < values; ++i) {
if (!std::isfinite(cpu_hdr[i]) || !std::isfinite(gpu_hdr[i])) return 1;
const double absolute = std::fabs(cpu_hdr[i] - gpu_hdr[i]);
max_abs = std::fmax(max_abs, absolute);
if (std::fabs(cpu_hdr[i]) > 1e-30)
@@ -66,5 +97,5 @@ int main() {
std::printf("HIP PSF cache comparison: max_abs=%.17g max_rel=%.17g\n", max_abs, max_rel);
hip_psf_sink_destroy(sink);
std::free(cpu_hdr); std::free(gpu_hdr); psf_kernel_cache_destroy(&cache);
return max_abs <= 1e-12 && max_rel <= 1e-12 ? 0 : 1;
return max_abs <= 1e-10 && max_rel <= 1e-12 ? 0 : 1;
}
tests/test_hip_renderer.py
+61
View File
@@ -0,0 +1,61 @@
#!/usr/bin/env python3
"""Sequential, bounded CPU/HIP renderer checks for shared-HDR fallback safety."""
import argparse
import importlib.util
import math
import os
from pathlib import Path
import re
import struct
import subprocess
import tempfile
def main():
parser = argparse.ArgumentParser(description=__doc__)
parser.add_argument("cpu", type=Path, help="Minkowski CPU executable")
parser.add_argument("hip", type=Path, help="Minkowski HIP executable")
args = parser.parse_args()
repo = Path(__file__).resolve().parents[1]
spec = importlib.util.spec_from_file_location("fits", repo / "scripts/fits_floatdiff.py")
fits = importlib.util.module_from_spec(spec)
spec.loader.exec_module(fits)
env = dict(os.environ, OMP_NUM_THREADS="4")
with tempfile.TemporaryDirectory(prefix="gr-hip-renderer-test-") as directory:
for min_y, expected in [("0", (2, 1, 0)), ("1e30", (0, 0, 3))]:
results = []
for label, binary in [("cpu", args.cpu), ("hip", args.hip)]:
output = Path(directory) / f"{label}-{min_y}.png"
command = [str(binary.resolve()), "--catalog",
str(repo / "tests/data/hip_psf_mixed_catalog.csv"),
"--width", "96", "--height", "64", "--fov-deg", "30",
"--look-ra-deg", "2", "--look-dec-deg", "1",
"--coarse-cell-pixels", "16", "--refine-max-level", "0",
"--exposure", "1000", "--psf-fwhm-pixels", "2.7",
"--psf-moffat-beta", "4.5", "--psf-relative-tail", "1e-8",
"--max-cache-psf-flux", "1", "--psf-min-y", min_y,
"--hdr-output", "--output", str(output)]
print("RUN", command, flush=True)
run = subprocess.run(command, env=env, text=True, capture_output=True, timeout=30)
print(run.stdout, run.stderr, flush=True)
run.check_returncode()
match = re.search(r"PSF splats: cached (\d+), cached wing-clipped (\d+), "
r"direct fallbacks (\d+), discarded below min-Y (\d+)",
run.stdout + run.stderr)
assert match, "missing splat accounting"
stats = tuple(map(int, match.groups()))
assert (stats[0], stats[2], stats[3]) == expected, stats
results.append((stats, fits.read_fits(output.with_name(output.stem + "_HDR.fits"))))
assert results[0][0] == results[1][0], "CPU/HIP statistics differ"
(shape, cpu), (gpu_shape, gpu) = results[0][1], results[1][1]
assert shape == gpu_shape
maximum = 0.0
for (a,), (b,) in zip(struct.iter_unpack(">f", cpu), struct.iter_unpack(">f", gpu)):
assert math.isfinite(a) and math.isfinite(b)
assert math.isclose(a, b, rel_tol=2**-23, abs_tol=0), (a, b)
maximum = max(maximum, abs(a - b))
print(f"min_y={min_y}: statistics match; HDR max_abs={maximum}; exact={cpu == gpu}", flush=True)
if __name__ == "__main__":
main()