Compare commits
2
Commits
| Author | SHA1 | Date | |
|---|---|---|---|
 wyj
|
ab7d879f6c | ||
|
wyj
|
e1ec480669 |
No files matched your search
@@ -41,9 +41,10 @@ SCHWARZSCHILD_TEST_TARGET := $(BUILD_DIR)/test_schwarzschild
|
|||||||
OBSERVER_TRACK_TEST_TARGET := $(BUILD_DIR)/test_observer_track
|
OBSERVER_TRACK_TEST_TARGET := $(BUILD_DIR)/test_observer_track
|
||||||
CATALOG_PREFETCH_TEST_TARGET := $(BUILD_DIR)/test_catalog_prefetch
|
CATALOG_PREFETCH_TEST_TARGET := $(BUILD_DIR)/test_catalog_prefetch
|
||||||
HIP_PSF_TEST_TARGET := $(BUILD_DIR)/test_hip_psf
|
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
|
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))
|
ifneq ($(filter 0 1,$(PSF_EVENT_SINK)),$(PSF_EVENT_SINK))
|
||||||
$(error Unknown PSF_EVENT_SINK '$(PSF_EVENT_SINK)'; choose 0 or 1)
|
$(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.
|
# It requires PSF_BACKEND=hip and an actual HIP GPU agent at execution time.
|
||||||
ifeq ($(PSF_BACKEND),hip)
|
ifeq ($(PSF_BACKEND),hip)
|
||||||
hip-psf-test: $(HIP_PSF_TEST_TARGET)
|
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)
|
$(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
|
else
|
||||||
hip-psf-test:
|
hip-psf-test hip-psf-bench:
|
||||||
$(error hip-psf-test requires PSF_BACKEND=hip)
|
$(error HIP PSF checks require PSF_BACKEND=hip)
|
||||||
endif
|
endif
|
||||||
|
|
||||||
run: $(TARGET)
|
run: $(TARGET)
|
||||||
|
|||||||
@@ -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
|
C with OpenMP CPU parallelism; an optional HIP backend accelerates PSF
|
||||||
accumulation.
|
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 [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;
|
The current scope is black-hole capture and distant stellar backgrounds;
|
||||||
local matter emission, accretion disks, and plasma are outside this stage.
|
local matter emission, accretion disks, and plasma are outside this stage.
|
||||||
|
|||||||
@@ -19,6 +19,8 @@
|
|||||||
|
|
||||||
当前实现支持解析 **Minkowski** 与 **Schwarzschild** 时空、单张图像、基于观测者轨迹的图像序列、自适应透镜网格，以及可复用的透镜映射文件。程序主要使用 C 编写，通过 OpenMP 实现 CPU 并行；可选的 HIP 后端用于加速 PSF 累积。
|
当前实现支持解析 **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)。
|
[Nmesh](https://github.com/nmeshsource/nmesh) 数值时空后端和双黑洞（BBH）渲染仍在规划中。当前范围为黑洞捕获与远处恒星背景，暂不包含局域物质辐射、吸积盘或等离子体。架构与开发路线参见[设计文档](nr_spacetime_movie_renderer_design.md)。
|
||||||
|
|
||||||
## 编译
|
## 编译
|
||||||
|
|||||||
@@ -1,3 +1,7 @@
|
|||||||
|
> 历史记录：此运行早于 `265d7b95d5b916ec4768d174d73ad963a43f0088` 的伪像修复。
|
||||||
|
> 当前 CPU/HIP 配对基准及优化提交 hash 见 [完整 4K 新记录](2mass_galactic_center_cpu_hip_2026-09-07.md)。
|
||||||
|
> 下方原始命令与输出保留，不用于新版本计数对照。
|
||||||
|
|
||||||
Git hash: dd8cf29dac557c881e5df1db15a58fae15294967
|
Git hash: dd8cf29dac557c881e5df1db15a58fae15294967
|
||||||
|
|
||||||
```
|
```
|
||||||
|
|||||||
@@ -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 加速比
|
||||||
|
替代这里的完整运行结果。
|
||||||
@@ -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
|
||||||
@@ -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
|
||||||
@@ -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
|
||||||
|
```
|
||||||
@@ -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
|
||||||
@@ -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
|
||||||
@@ -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
|
||||||
@@ -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
|
||||||
@@ -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
|
||||||
@@ -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
|
||||||
@@ -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
|
||||||
@@ -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'
|
||||||
@@ -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
|
||||||
@@ -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
|
||||||
@@ -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
|
||||||
@@ -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
|
||||||
|
File diff suppressed because it is too large.
Load diff
@@ -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
|
||||||
@@ -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
|
||||||
@@ -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
|
||||||
@@ -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
|
||||||
@@ -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
|
||||||
@@ -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
|
||||||
@@ -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
|
||||||
@@ -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
|
||||||
@@ -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)
|
||||||
@@ -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
|
||||||
@@ -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
|
||||||
@@ -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"
|
||||||
|
}
|
||||||
|
}
|
||||||
@@ -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())
|
||||||
@@ -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
|
||||||
@@ -0,0 +1 @@
|
|||||||
|
HIP PSF cache comparison: max_abs=2.8421709430404007e-14 max_rel=7.5319184269262976e-16
|
||||||
@@ -0,0 +1 @@
|
|||||||
|
HIP PSF cache comparison: max_abs=2.8421709430404007e-14 max_rel=7.5319184269262976e-16
|
||||||
@@ -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
|
||||||
@@ -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
|
||||||
@@ -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
|
||||||
@@ -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
|
reference path with ordered HDR transfers. HIP initialization or execution
|
||||||
errors terminate the render instead of silently selecting the CPU backend.
|
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
|
## Rebuilding and checks
|
||||||
|
|
||||||
Build type, spacetime, HDR, and PSF backend configurations have separate object
|
Build type, spacetime, HDR, and PSF backend configurations have separate object
|
||||||
|
|||||||
@@ -1,5 +1,20 @@
|
|||||||
# GPU PSF 累积计划
|
# 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. 目的与边界
|
## 1. 目的与边界
|
||||||
|
|
||||||
本文件规划的是**最终点源 PSF 到 HDR framebuffer 的累积**加速，而不是 GPU
|
本文件规划的是**最终点源 PSF 到 HDR framebuffer 的累积**加速，而不是 GPU
|
||||||
|
|||||||
@@ -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
|
# 23. Metric evaluator 的 thread-local workspace
|
||||||
|
|
||||||
nmesh 单点查询需要：
|
nmesh 单点查询需要：
|
||||||
|
|||||||
+133
-22
@@ -133,6 +133,9 @@ typedef struct {
|
|||||||
const PsfKernelCache *cache;
|
const PsfKernelCache *cache;
|
||||||
#ifdef PSF_BACKEND_HIP
|
#ifdef PSF_BACKEND_HIP
|
||||||
HipPsfSink *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];
|
char hip_message[256];
|
||||||
HipPsfTiming hip_timing;
|
HipPsfTiming hip_timing;
|
||||||
#endif
|
#endif
|
||||||
@@ -176,7 +179,7 @@ static int psf_event_sink_init(PsfEventSink *sink, double *hdr, int width,
|
|||||||
return 0;
|
return 0;
|
||||||
}
|
}
|
||||||
|
|
||||||
static int psf_event_sink_flush(PsfEventSink *sink) {
|
static int psf_event_sink_flush_unlocked(PsfEventSink *sink) {
|
||||||
if (sink->failed)
|
if (sink->failed)
|
||||||
return -1;
|
return -1;
|
||||||
#ifdef PSF_BACKEND_HIP
|
#ifdef PSF_BACKEND_HIP
|
||||||
@@ -197,8 +200,27 @@ static int psf_event_sink_flush(PsfEventSink *sink) {
|
|||||||
return 0;
|
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) {
|
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))
|
if (psf_event_sink_flush(sink))
|
||||||
|
#endif
|
||||||
return -1;
|
return -1;
|
||||||
#ifdef PSF_BACKEND_HIP
|
#ifdef PSF_BACKEND_HIP
|
||||||
if (sink->hip != NULL && hip_psf_sink_finish(sink->hip, sink->hdr,
|
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) {
|
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);
|
int result = psf_event_sink_finish_for_cpu(sink);
|
||||||
#ifdef PSF_BACKEND_HIP
|
#ifdef PSF_BACKEND_HIP
|
||||||
if (sink->hip != NULL && hip_psf_sink_get_timing(sink->hip, &sink->hip_timing)) {
|
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;
|
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) {
|
static void psf_event_sink_emit(PsfEventSink *sink, const PsfCachedEvent *event) {
|
||||||
if (sink->failed)
|
if (sink->failed)
|
||||||
return;
|
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,
|
&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);
|
context->max_cache_psf_flux, context->psf_relative_tail, context->psf_min_y);
|
||||||
if (direct_fallback == 1) {
|
if (direct_fallback == 1) {
|
||||||
/* A direct fallback forms an explicit ordered CPU/GPU boundary. */
|
/* Exclude every other submit and fallback until HDR is back on device. */
|
||||||
if (psf_event_sink_finish_for_cpu(context->event_sink))
|
#ifdef PSF_BACKEND_HIP
|
||||||
return -1;
|
const double submit_start = omp_get_wtime();
|
||||||
splat_moffat_direct(context->hdr, context->width, context->height, image_x,
|
if (context->event_sink->hip_lock)
|
||||||
image_y, color, flux, context->psf,
|
omp_set_lock(context->event_sink->hip_lock);
|
||||||
context->psf_relative_tail, context->psf_min_y);
|
#endif
|
||||||
if (psf_event_sink_resume_gpu(context->event_sink))
|
int failed = psf_event_sink_finish_for_cpu(context->event_sink);
|
||||||
return -1;
|
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) {
|
} 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);
|
psf_event_sink_emit(context->event_sink, &event);
|
||||||
#else
|
#else
|
||||||
splat_prepared_cached_event(context->hdr, context->width, context->height,
|
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);
|
mesh->triangle_count);
|
||||||
|
|
||||||
#ifdef PSF_BACKEND_HIP
|
#ifdef PSF_BACKEND_HIP
|
||||||
/* A single GPU HDR framebuffer owns all cached events. Keep catalog/lens
|
/* CPU workers own only bounded event chunks. A single device cache/HDR is
|
||||||
* work serial for this first direct-atomic integration; the CPU parallel
|
* borrowed under a coarse submission lock, never duplicated per worker. */
|
||||||
* private-HDR path remains the PSF_BACKEND=cpu implementation. */
|
PsfEventSink owner;
|
||||||
const CatalogSplatStats hip_stats = splat_catalog_triangles(
|
if (psf_event_sink_init(&owner, hdr, width, height, psf_cache)) return SIZE_MAX;
|
||||||
mesh, catalog, hdr, width, height, exposure, psf, psf_cache,
|
omp_lock_t submit_lock;
|
||||||
max_magnification, max_cache_psf_flux, psf_relative_tail, psf_min_y,
|
omp_init_lock(&submit_lock);
|
||||||
0, mesh->triangle_count, NULL);
|
size_t hip_images = 0, hip_direct = 0, hip_clipped = 0, hip_discarded = 0;
|
||||||
copy_psf_splat_stats(psf_stats, hip_stats);
|
int hip_failed = 0, hip_workers = 0;
|
||||||
if (hip_stats.failed)
|
double hip_generate_seconds = 0.0, hip_submit_seconds = 0.0;
|
||||||
return SIZE_MAX;
|
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)
|
if (progress != NULL && progress->callback != NULL)
|
||||||
progress->callback(progress->context, FRAME_SPLAT_PROGRESS_END,
|
progress->callback(progress->context, FRAME_SPLAT_PROGRESS_END,
|
||||||
mesh->triangle_count, mesh->triangle_count);
|
mesh->triangle_count, mesh->triangle_count);
|
||||||
return hip_stats.images;
|
return hip_images;
|
||||||
#endif
|
#endif
|
||||||
|
|
||||||
const size_t pixel_count = (size_t)width * height * 3;
|
const size_t pixel_count = (size_t)width * height * 3;
|
||||||
|
|||||||
+6
-4
@@ -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,
|
int hip_psf_sink_load_hdr(HipPsfSink *sink, const double *hdr,
|
||||||
char *message, size_t message_size);
|
char *message, size_t message_size);
|
||||||
|
|
||||||
/* Queues a cached-event chunk. Direct fallbacks and min-Y discards stay with
|
/* Copies a chunk into one of two owned pinned staging slots; caller memory
|
||||||
* the caller; event_count must not exceed the creation capacity. */
|
* 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,
|
int hip_psf_sink_submit(HipPsfSink *sink, const PsfCachedEvent *events,
|
||||||
size_t event_count, char *message, size_t message_size);
|
size_t event_count, char *message, size_t message_size);
|
||||||
|
|
||||||
/* Completes all work and overwrites hdr with double linear RGB device HDR. */
|
/* Completes all work and overwrites hdr with double linear RGB device HDR. */
|
||||||
int hip_psf_sink_finish(HipPsfSink *sink, double *hdr, char *message,
|
int hip_psf_sink_finish(HipPsfSink *sink, double *hdr, char *message,
|
||||||
size_t message_size);
|
size_t message_size);
|
||||||
/* Valid after hip_psf_sink_finish(); timings are accumulated without a
|
/* Complete after finish(); otherwise includes recycled/completed slots only.
|
||||||
* per-batch host synchronization. */
|
* Every batch is timed with a fixed-size reusable event pool. */
|
||||||
int hip_psf_sink_get_timing(const HipPsfSink *sink, HipPsfTiming *timing);
|
int hip_psf_sink_get_timing(const HipPsfSink *sink, HipPsfTiming *timing);
|
||||||
void hip_psf_sink_destroy(HipPsfSink *sink);
|
void hip_psf_sink_destroy(HipPsfSink *sink);
|
||||||
|
|
||||||
|
|||||||
+103
-64
@@ -5,27 +5,34 @@
|
|||||||
#include <cmath>
|
#include <cmath>
|
||||||
#include <cstdio>
|
#include <cstdio>
|
||||||
#include <limits>
|
#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 upload_start = nullptr, upload_end = nullptr;
|
||||||
hipEvent_t kernel_start = nullptr, kernel_end = nullptr;
|
hipEvent_t kernel_start = nullptr, kernel_end = nullptr;
|
||||||
|
size_t count = 0;
|
||||||
};
|
};
|
||||||
|
|
||||||
enum { HIP_PSF_MAX_TIMED_BATCHES = 64 };
|
|
||||||
|
|
||||||
struct HipPsfSink {
|
struct HipPsfSink {
|
||||||
PsfCachedEvent *events = nullptr;
|
HipPsfBatch slots[2];
|
||||||
float *weights = nullptr;
|
float *weights = nullptr;
|
||||||
double *hdr = nullptr;
|
double *hdr = nullptr;
|
||||||
hipStream_t stream = nullptr;
|
hipStream_t stream = nullptr;
|
||||||
size_t event_capacity = 0, hdr_values = 0;
|
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;
|
int width = 0, height = 0, phase_resolution = 0, radius_pixels = 0;
|
||||||
double max_radius_pixels = 0.0;
|
double max_radius_pixels = 0.0;
|
||||||
hipEvent_t download_start = nullptr, download_end = nullptr;
|
hipEvent_t download_start = nullptr, download_end = nullptr;
|
||||||
std::vector<HipPsfBatchTiming> batches;
|
|
||||||
size_t measured_batches = 0;
|
|
||||||
HipPsfTiming timing = {};
|
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) {
|
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");
|
if (message && message_size) std::snprintf(message, message_size, "ok");
|
||||||
}
|
}
|
||||||
|
|
||||||
static void destroy_batch_timing(HipPsfBatchTiming *timing) {
|
static int collect_batch(HipPsfSink *sink, HipPsfBatch *batch,
|
||||||
if (!timing) return;
|
char *message, size_t message_size) {
|
||||||
(void)hipEventDestroy(timing->upload_start);
|
if (!batch->count) return 0;
|
||||||
(void)hipEventDestroy(timing->upload_end);
|
float upload_ms = 0.0f, kernel_ms = 0.0f;
|
||||||
(void)hipEventDestroy(timing->kernel_start);
|
if (report(hipEventSynchronize(batch->kernel_end), message, message_size) ||
|
||||||
(void)hipEventDestroy(timing->kernel_end);
|
report(hipEventElapsedTime(&upload_ms, batch->upload_start, batch->upload_end), message, message_size) ||
|
||||||
*timing = {};
|
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,
|
__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 width, int height, const float *weights,
|
||||||
int phase_resolution, int radius_pixels,
|
int phase_resolution, int radius_pixels,
|
||||||
double max_radius_pixels, double *hdr) {
|
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;
|
if (index >= count) return;
|
||||||
const PsfCachedEvent event = events[index];
|
const PsfCachedEvent event = events[index];
|
||||||
const int base_x = (int)floor(event.x), base_y = (int)floor(event.y);
|
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 (!row_range(event.support_radius, fx, fy, support, offset_y, &first, &last)) continue;
|
||||||
if (first < -base_x) first = -base_x;
|
if (first < -base_x) first = -base_x;
|
||||||
if (last >= width - base_x) last = width - base_x - 1;
|
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 int px = base_x + offset_x;
|
||||||
const double w00 = weights[weight_index(phase_resolution, radius_pixels, x0, y0, offset_x, offset_y)];
|
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)];
|
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,
|
const PsfKernelCache *cache, size_t event_capacity,
|
||||||
char *message, size_t message_size) {
|
char *message, size_t message_size) {
|
||||||
if (!out || width <= 0 || height <= 0 || !cache || !cache->ready || !cache->weights ||
|
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");
|
if (message && message_size) std::snprintf(message, message_size, "invalid HIP PSF sink arguments");
|
||||||
return -1;
|
return -1;
|
||||||
}
|
}
|
||||||
const size_t nodes = (size_t)cache->phase_resolution + 1;
|
const size_t nodes = (size_t)cache->phase_resolution + 1;
|
||||||
const size_t side = (size_t)cache->radius_pixels * 2 + 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 ||
|
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 ||
|
nodes * nodes * side > std::numeric_limits<size_t>::max() / side ||
|
||||||
(size_t)width * (size_t)height > std::numeric_limits<size_t>::max() / 3) {
|
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");
|
if (message && message_size) std::snprintf(message, message_size, "HIP PSF sink size overflow");
|
||||||
return -1;
|
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->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;
|
sink->radius_pixels = cache->radius_pixels; sink->max_radius_pixels = cache->max_radius_pixels;
|
||||||
const size_t weight_count = nodes * nodes * side * side;
|
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_start), message, message_size) ||
|
||||||
report(hipEventCreate(&sink->download_end), 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->weights, weight_count * sizeof *sink->weights), message, message_size) ||
|
||||||
report(hipMalloc(&sink->hdr, sink->hdr_values * sizeof *sink->hdr), 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(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)) {
|
report(hipMemsetAsync(sink->hdr, 0, sink->hdr_values * sizeof *sink->hdr, sink->stream), message, message_size)) {
|
||||||
hip_psf_sink_destroy(sink); return -1;
|
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;
|
*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;
|
return -1;
|
||||||
}
|
}
|
||||||
if (!event_count) { ok(message, message_size); return 0; }
|
if (!event_count) { ok(message, message_size); return 0; }
|
||||||
const unsigned int threads = 128;
|
if (report(hipSetDevice(sink->device), message, message_size)) return -1;
|
||||||
const size_t blocks = (event_count + threads - 1) / threads;
|
const unsigned int threads = PSF_THREADS_PER_BLOCK;
|
||||||
HipPsfBatchTiming timing;
|
const size_t per_block = PSF_THREADS_PER_BLOCK / PSF_LANES_PER_EVENT;
|
||||||
const int timed = sink->batches.size() < HIP_PSF_MAX_TIMED_BATCHES;
|
const size_t blocks = event_count / per_block + (event_count % per_block != 0);
|
||||||
if (blocks > std::numeric_limits<unsigned int>::max() ||
|
if (blocks > std::numeric_limits<unsigned int>::max()) {
|
||||||
(timed && (report(hipEventCreate(&timing.upload_start), message, message_size) ||
|
if (message && message_size) std::snprintf(message, message_size, "HIP PSF grid too large");
|
||||||
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);
|
|
||||||
return -1;
|
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,
|
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);
|
sink->phase_resolution, sink->radius_pixels, sink->max_radius_pixels, sink->hdr);
|
||||||
if (report(hipGetLastError(), message, message_size) ||
|
if (report(hipGetLastError(), message, message_size) ||
|
||||||
(timed && report(hipEventRecord(timing.kernel_end, sink->stream), message, message_size))) {
|
report(hipEventRecord(batch.kernel_end, sink->stream), message, message_size)) return -1;
|
||||||
destroy_batch_timing(&timing);
|
batch.count = event_count;
|
||||||
return -1;
|
|
||||||
}
|
|
||||||
if (timed) {
|
|
||||||
sink->batches.push_back(timing);
|
|
||||||
++sink->timing.timed_batch_count;
|
|
||||||
}
|
|
||||||
++sink->timing.batch_count;
|
++sink->timing.batch_count;
|
||||||
sink->timing.event_count += event_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;
|
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");
|
if (message && message_size) std::snprintf(message, message_size, "invalid HIP PSF HDR upload");
|
||||||
return -1;
|
return -1;
|
||||||
}
|
}
|
||||||
if (report(hipMemcpyAsync(sink->hdr, hdr, sink->hdr_values * sizeof *hdr,
|
if (report(hipSetDevice(sink->device), message, message_size) ||
|
||||||
hipMemcpyHostToDevice, sink->stream), 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;
|
return -1;
|
||||||
ok(message, message_size); return 0;
|
ok(message, message_size); return 0;
|
||||||
}
|
}
|
||||||
|
|
||||||
extern "C" int hip_psf_sink_finish(HipPsfSink *sink, double *hdr, char *message, size_t message_size) {
|
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 (!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(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(hipEventRecord(sink->download_end, sink->stream), message, message_size) ||
|
||||||
report(hipStreamSynchronize(sink->stream), message, message_size)) return -1;
|
report(hipStreamSynchronize(sink->stream), message, message_size)) return -1;
|
||||||
float milliseconds = 0.0f;
|
float milliseconds = 0.0f;
|
||||||
for (size_t i = sink->measured_batches; i < sink->batches.size(); ++i) {
|
for (HipPsfBatch &batch : sink->slots)
|
||||||
const HipPsfBatchTiming &batch = sink->batches[i];
|
if (collect_batch(sink, &batch, message, message_size)) return -1;
|
||||||
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();
|
|
||||||
if (report(hipEventElapsedTime(&milliseconds, sink->download_start, sink->download_end),
|
if (report(hipEventElapsedTime(&milliseconds, sink->download_start, sink->download_end),
|
||||||
message, message_size)) return -1;
|
message, message_size)) return -1;
|
||||||
sink->timing.download_seconds += milliseconds * 1e-3;
|
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) {
|
extern "C" void hip_psf_sink_destroy(HipPsfSink *sink) {
|
||||||
if (!sink) return;
|
if (!sink) return;
|
||||||
for (HipPsfBatchTiming &timing : sink->batches)
|
(void)hipSetDevice(sink->device);
|
||||||
destroy_batch_timing(&timing);
|
if (sink->stream) (void)hipStreamSynchronize(sink->stream);
|
||||||
(void)hipEventDestroy(sink->download_start);
|
for (HipPsfBatch &batch : sink->slots) {
|
||||||
(void)hipEventDestroy(sink->download_end);
|
if (batch.upload_start) (void)hipEventDestroy(batch.upload_start);
|
||||||
(void)hipFree(sink->events); (void)hipFree(sink->weights); (void)hipFree(sink->hdr);
|
if (batch.upload_end) (void)hipEventDestroy(batch.upload_end);
|
||||||
(void)hipStreamDestroy(sink->stream); delete sink;
|
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;
|
||||||
}
|
}
|
||||||
@@ -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;
|
||||||
|
}
|
||||||
+32
-1
@@ -57,7 +57,38 @@ int main() {
|
|||||||
return 1;
|
return 1;
|
||||||
}
|
}
|
||||||
double max_abs = 0.0, max_rel = 0.0;
|
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) {
|
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]);
|
const double absolute = std::fabs(cpu_hdr[i] - gpu_hdr[i]);
|
||||||
max_abs = std::fmax(max_abs, absolute);
|
max_abs = std::fmax(max_abs, absolute);
|
||||||
if (std::fabs(cpu_hdr[i]) > 1e-30)
|
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);
|
std::printf("HIP PSF cache comparison: max_abs=%.17g max_rel=%.17g\n", max_abs, max_rel);
|
||||||
hip_psf_sink_destroy(sink);
|
hip_psf_sink_destroy(sink);
|
||||||
std::free(cpu_hdr); std::free(gpu_hdr); psf_kernel_cache_destroy(&cache);
|
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;
|
||||||
}
|
}
|
||||||
@@ -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()
|
||||||
Reference in new issue
Block a user