feat(ugc): ray_backend=hiprt traces on the GPU with HIPRT and Orochi (optional build)

HIPRT (MIT) behind the CMake option DLU_HIPRT (off): its headers come from its SDK (HIPRT_ROOT, else ROCm's /opt/rocm) and are copied next to the servers; its library is loaded when first used (hiprtew), as HIP or CUDA are by Orochi (MIT, fetched pinned by hash; CUDA when its toolkit is found). The trace kernels (nearest hit skipping the triangle a ray leaves, any hit) are compiled the first time and kept in cache/hiprt. One GPU context for the process (hiprt_device picks the GPU); the workers take turns on it. When HIPRT, the GPU or a scene's upload fails, Embree is used instead, and the UGC server logs why at start. For a GPU the rays go in batches (UgcRays::Scene gets batch queries; the CPU backends answer them a ray at a time): - hidden faces: with a batch backend the paths are traced side by side, a bounce at a time (the path code split into Start, Scatter and Bounce, the one by one tracing unchanged); the same paths with the same random numbers, so the same triangles are decided (tested with builtin side by side) - the occlusion bake and the denoised icons' traced occlusion always ask in batches (the same rays, the same results) Its symbols are hidden: the servers export theirs (-rdynamic), and HIPRT's library, which has an Orochi of its own, would otherwise call ours. Check: configure with -DDLU_HIPRT=ON on a machine with ROCm (or HIPRT's SDK) and an AMD RDNA or NVIDIA GPU; UgcServer --make-model x.lxfml out hiprt; the UGC tests (hits, hidden faces and occlusion against builtin); a build without it leaves everything as before. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
2026-10-05 12:23:44 +00:00 · 2026-09-29 11:57:37 -05:00
parent 2e2e8153e2
commit 07dbf30dc8
13 changed files with 719 additions and 92 deletions
--- a/dDashboardServer/routes/SettingsCatalog.cpp
+++ b/dDashboardServer/routes/SettingsCatalog.cpp
@@ -504,6 +504,7 @@ namespace {
 		c.Add(Labels(Choice(UGC, "hsr_method", "How hidden faces are found", "toolbox: LU Toolbox's paths (above). fast: the model rendered from 42 directions around it, removing what shows in none; much faster, but it also removes faces seen only by bounced light (insides seen through openings, recesses). Staff can pick another for one make when making models again.", "toolbox", { "toolbox", "fast" }), { "LU Toolbox (paths)", "Fast (renders from around)" }));
 		c.Add(Unit(Int(UGC, "hsr_fast_resolution", "Fast method's render size", "The size of each of the fast method's 42 renders; bigger keeps smaller visible faces.", "1024", 64, 4096), "pixels"));
 		c.Add(Labels(Choice(UGC, "ray_backend", "Ray tracer", "What traces the rays of the hidden faces' paths and of the occlusion: builtin (the UGC server's own), embree (Intel Embree on the CPU) or hiprt (the GPU, when the server was built with it and has one; else embree). They give the same results but for rounding. Staff can pick another for one make when making models again.", "builtin", { "builtin", "embree", "hiprt" }), { "Built in", "Embree (CPU)", "HIPRT (GPU)" }));
+		c.Add(Int(UGC, "hiprt_device", "GPU for HIPRT", "Which GPU the hiprt ray tracer uses: 0 is the first HIP (AMD) or CUDA (NVIDIA) device.", "0", 0, 16, true));
 		c.Add(Labels(Choice(UGC, "denoise", "Denoise icons", "oidn (when the server was built with Intel Open Image Denoise; else off): a model's icon is drawn from its colors before the occlusion bake, with its occlusion traced per pixel with a few rays and the noise removed by the denoiser, instead of the baked occlusion. The model itself keeps its baked occlusion (a denoiser only works on images).", "off", { "off", "oidn" }), { "Off", "Open Image Denoise" }));
 		c.Add(Int(UGC, "denoise_samples", "Denoised occlusion rays per pixel", "Occlusion rays traced from each pixel of a denoised icon (before it is scaled down, so 16 times as many per icon pixel); fewer is faster and noisier.", "4", 1, 256));
 		c.Add(Bool(UGC, "bake_ao", "Darken hidden corners", "Ambient occlusion baked into the vertex colors.", true));
--- a/dUgcServer/CMakeLists.txt
+++ b/dUgcServer/CMakeLists.txt
@@ -33,6 +33,18 @@ if(DLU_OIDN)
 	target_link_libraries(dUgc PRIVATE OpenImageDenoise)
 	target_compile_definitions(dUgc PRIVATE DLU_OIDN)
 endif()
+# HIPRT and Orochi, when built with DLU_HIPRT (thirdparty/CMakeLists.txt): the ray_backend=hiprt setting
+if(DLU_HIPRT)
+	target_sources(dUgc PRIVATE "Render/UgcRaysHiprt.cpp")
+	# hiprtew's function pointers have the names of HIPRT's functions: hidden, or HIPRT's library would call them
+	# instead of its own (the servers export their symbols)
+	if(NOT MSVC)
+		set_source_files_properties("Render/UgcRaysHiprt.cpp" PROPERTIES COMPILE_OPTIONS "-fvisibility=hidden")
+	endif()
+	target_include_directories(dUgc PRIVATE "${DLU_HIPRT_INCLUDE_DIR}")
+	target_link_libraries(dUgc PRIVATE dlu_orochi)
+	target_compile_definitions(dUgc PRIVATE DLU_HIPRT DLU_HIPRT_INCLUDE_DIR="${DLU_HIPRT_INCLUDE_DIR}")
+endif()

 add_executable(UgcServer "UgcServer.cpp" "Processing/UgcProcessor.cpp" "Bricks/UgcCdClient.cpp")

--- a/dUgcServer/Render/UgcHsr.cpp
+++ b/dUgcServer/Render/UgcHsr.cpp
@@ -343,6 +343,79 @@ namespace {
 			return m_Mesh.positions[m_Mesh.indices[t * 3]] * w.x + m_Mesh.positions[m_Mesh.indices[t * 3 + 1]] * w.y + m_Mesh.positions[m_Mesh.indices[t * 3 + 2]] * w.z;
 		}

+		// A path between its rays
+		struct Path {
+			glm::vec3 ng{}, n{}, incoming{}, p{};
+			glm::vec3 direction{}; // of the ray it waits for
+			uint32_t self{};       // the triangle it is on (the ray never hits it)
+			float throughput{ 1.0f };
+			float minRayPdf{ INF };
+			int glossy{};
+			int bounce{};
+			bool sampleGlossy{};
+			Random random;
+		};
+
+		enum class eStep : uint8_t { RAY, ESCAPED, ENDED };
+
+		// A path from point w (barycentric) of triangle t, its random numbers from `seed`
+		Path Start(uint32_t t, const glm::vec3& w, uint64_t seed) const {
+			Path path{ .random = Random(seed) };
+			path.ng = FaceNormal(t);
+			// The bake looks at the point along its smooth normal (the ray comes from there): that side is lit, and
+			// when the winding faces the other way Cycles treats it as the back of the face (both normals turned)
+			path.n = ShadingNormal(t, w, path.ng);
+			path.incoming = path.n;
+			if (glm::dot(path.ng, path.incoming) < 0.0f) {
+				path.ng = -path.ng;
+				path.n = -path.n;
+			}
+			path.p = Position(t, w);
+			path.self = t;
+			return path;
+		}
+
+		// The direction the path goes on in: RAY (from path.p along path.direction), or ENDED
+		eStep Scatter(Path& path) const {
+			// The material's normal: the Principled BSDF keeps its mirror reflection above the surface
+			const Principled material(EnsureValidReflection(path.ng, path.incoming, path.n), path.incoming, path.ng, path.n, path.bounce == 0, path.minRayPdf);
+			const auto sample = material.Draw(path.random);
+			// Nothing sampled (below the surface): the path ends
+			if (!(sample.pdf > 0.0f) || !(sample.throughput > 0.0f)) return eStep::ENDED;
+			path.direction = sample.direction;
+			path.throughput *= sample.throughput;
+			path.minRayPdf = std::min(path.minRayPdf, sample.pdf);
+			path.sampleGlossy = sample.glossy;
+			return eStep::RAY;
+		}
+
+		// What the ray hit: ESCAPED (nothing: the sky), ENDED, or RAY (it bounced: Scatter again)
+		eStep Bounce(Path& path, const UgcRays::Hit& hit) const {
+			const auto& direction = path.direction;
+			if (m_Options.groundPlane && GroundHit(path.p, direction, hit.t) < hit.t) return eStep::ENDED;
+			if (hit.triangle == UgcRays::NONE) return eStep::ESCAPED;
+			// Past the bounce limits the next surface doesn't scatter (Cycles: Max Bounces, Glossy 4)
+			if (path.bounce + 1 > m_Options.bounces) return eStep::ENDED;
+			if (path.sampleGlossy && ++path.glossy > GLOSSY_BOUNCES) return eStep::ENDED;
+			if (path.bounce + 1 >= 2) {
+				const float probability = std::min(std::sqrt(path.throughput), 1.0f);
+				if (path.random.Float() >= probability) return eStep::ENDED;
+				path.throughput /= probability;
+			}
+			path.self = hit.triangle;
+			path.ng = FaceNormal(path.self);
+			path.n = ShadingNormal(path.self, glm::vec3(1.0f - hit.u - hit.v, hit.u, hit.v), path.ng);
+			path.incoming = -direction;
+			// Hit from behind: the back is lit (both normals turned towards the ray)
+			if (glm::dot(path.ng, path.incoming) < 0.0f) {
+				path.ng = -path.ng;
+				path.n = -path.n;
+			}
+			path.p = path.p + direction * hit.t;
+			path.bounce++;
+			return eStep::RAY;
+		}
+
 		/**
 		 * One path from point w (barycentric) of triangle t, as a Cycles diffuse bake traces it: whether it reaches the
 		 * sky. The path bounces off the model in the directions the bake material draws until a bounce ray hits
@@ -351,54 +424,17 @@ namespace {
 		 * second bounce on the path may end early (Cycles' Russian roulette: it goes on with probability
 		 * sqrt(throughput)).
 		 */
-		bool Escapes(uint32_t t, const glm::vec3& w, Random& random) const {
-			glm::vec3 ng = FaceNormal(t);
-			// The bake looks at the point along its smooth normal (the ray comes from there): that side is lit, and
-			// when the winding faces the other way Cycles treats it as the back of the face (both normals turned)
-			glm::vec3 n = ShadingNormal(t, w, ng);
-			glm::vec3 incoming = n;
-			if (glm::dot(ng, incoming) < 0.0f) {
-				ng = -ng;
-				n = -n;
-			}
-			glm::vec3 p = Position(t, w);
-			uint32_t self = t;
-			float throughput = 1.0f;
-			float minRayPdf = INF;
-			int glossy = 0;
-			for (int bounce = 0;; bounce++) {
-				// The material's normal: the Principled BSDF keeps its mirror reflection above the surface
-				const Principled material(EnsureValidReflection(ng, incoming, n), incoming, ng, n, bounce == 0, minRayPdf);
-				const auto sample = material.Draw(random);
-				// Nothing sampled (below the surface): the path ends
-				if (!(sample.pdf > 0.0f) || !(sample.throughput > 0.0f)) return false;
-				const auto& direction = sample.direction;
-				throughput *= sample.throughput;
-				minRayPdf = std::min(minRayPdf, sample.pdf);
-				const auto hit = m_Rays->Closest(p, direction, self);
-				if (m_Options.groundPlane && GroundHit(p, direction, hit.t) < hit.t) return false;
-				if (hit.triangle == UgcRays::NONE) return true;
-				// Past the bounce limits the next surface doesn't scatter (Cycles: Max Bounces, Glossy 4)
-				if (bounce + 1 > m_Options.bounces) return false;
-				if (sample.glossy && ++glossy > GLOSSY_BOUNCES) return false;
-				if (bounce + 1 >= 2) {
-					const float probability = std::min(std::sqrt(throughput), 1.0f);
-					if (random.Float() >= probability) return false;
-					throughput /= probability;
-				}
-				self = hit.triangle;
-				ng = FaceNormal(self);
-				n = ShadingNormal(self, glm::vec3(1.0f - hit.u - hit.v, hit.u, hit.v), ng);
-				incoming = -direction;
-				// Hit from behind: the back is lit (both normals turned towards the ray)
-				if (glm::dot(ng, incoming) < 0.0f) {
-					ng = -ng;
-					n = -n;
-				}
-				p = p + direction * hit.t;
+		bool Escapes(uint32_t t, const glm::vec3& w, uint64_t seed) const {
+			auto path = Start(t, w, seed);
+			while (Scatter(path) == eStep::RAY) {
+				const auto step = Bounce(path, m_Rays->Closest(path.p, path.direction, path.self));
+				if (step != eStep::RAY) return step == eStep::ESCAPED;
 			}
+			return false;
 		}

+		const UgcRays::Scene& Rays() const { return *m_Rays; }
+
 	private:
 		const UgcModel::Mesh& m_Mesh;
 		std::unique_ptr<UgcRays::Scene> m_Rays; // the nearest hit (never the triangle a ray leaves, as in Cycles)
@@ -473,6 +509,77 @@ namespace UgcHsr {
 		const Tracer tracer(mesh, options);
 		const int samples = std::max(options.samples, 1);
 		uint64_t points = 0, paths = 0, sinceCheckpoint = 0;
+		if (tracer.Rays().PrefersBatches() || options.sideBySide) {
+			// Side by side (a GPU): the triangles in groups of about GROUP_PATHS paths; in each round every point of every
+			// triangle of the group not seen yet starts a path, and the round's paths are traced together a bounce at a
+			// time. Each path is the one traced one by one below (the same random numbers), and a triangle is kept when
+			// any of its paths escapes, so the triangles decided are the same; only paths the one by one tracing would
+			// have skipped after one escaped are traced too.
+			constexpr size_t GROUP_PATHS = 1u << 18;
+			uint32_t next = 0;
+			while (next < triangles) {
+				std::vector<uint32_t> group;
+				std::vector<std::vector<glm::vec3>> groupWeights;
+				size_t groupPaths = 0;
+				for (; next < triangles && groupPaths < GROUP_PATHS; next++) {
+					// A triangle without area draws nothing (LU Toolbox's bake leaves it dark too): removed
+					if (tracer.FaceNormal(next) == glm::vec3(0.0f)) {
+						visible[next] = false;
+						continue;
+					}
+					const auto& a = mesh.positions[mesh.indices[next * 3]];
+					const auto& b = mesh.positions[mesh.indices[next * 3 + 1]];
+					const auto& c = mesh.positions[mesh.indices[next * 3 + 2]];
+					groupWeights.push_back(SamplePoints(a, b, c, options.spacing, static_cast<size_t>(std::max(options.minPoints, 1))));
+					group.push_back(next);
+					groupPaths += groupWeights.back().size();
+					points += groupWeights.back().size();
+				}
+				std::vector<uint8_t> escaped(group.size(), 0);
+				for (int sample = 0; sample < samples; sample++) {
+					std::vector<Tracer::Path> live;
+					std::vector<uint32_t> owner; // the path's triangle in the group
+					for (uint32_t g = 0; g < group.size(); g++) {
+						if (escaped[g]) continue;
+						for (size_t i = 0; i < groupWeights[g].size(); i++) {
+							auto path = tracer.Start(group[g], groupWeights[g][i], PathSeed(options.seed, group[g], i, static_cast<uint64_t>(sample)));
+							paths++;
+							if (tracer.Scatter(path) != Tracer::eStep::RAY) continue;
+							live.push_back(std::move(path));
+							owner.push_back(g);
+						}
+					}
+					if (live.empty()) break;
+					std::vector<UgcRays::Ray> rays;
+					std::vector<UgcRays::Hit> hits;
+					while (!live.empty()) {
+						UgcThrottle::Checkpoint();
+						rays.resize(live.size());
+						hits.resize(live.size());
+						for (size_t k = 0; k < live.size(); k++) rays[k] = UgcRays::Ray{ live[k].p, 0.0f, live[k].direction, INF, live[k].self };
+						tracer.Rays().Closest(rays.data(), hits.data(), live.size());
+						size_t kept = 0;
+						for (size_t k = 0; k < live.size(); k++) {
+							if (escaped[owner[k]]) continue;
+							const auto step = tracer.Bounce(live[k], hits[k]);
+							if (step == Tracer::eStep::ESCAPED) escaped[owner[k]] = 1;
+							if (step != Tracer::eStep::RAY || tracer.Scatter(live[k]) != Tracer::eStep::RAY) continue;
+							if (kept != k) {
+								live[kept] = std::move(live[k]);
+								owner[kept] = owner[k];
+							}
+							kept++;
+						}
+						live.erase(live.begin() + static_cast<std::ptrdiff_t>(kept), live.end());
+						owner.resize(kept);
+					}
+				}
+				for (size_t g = 0; g < group.size(); g++) visible[group[g]] = escaped[g] != 0;
+			}
+			if (pointCount) *pointCount = points;
+			if (pathCount) *pathCount = paths;
+			return visible;
+		}
 		for (uint32_t t = 0; t < triangles; t++) {
 			const auto& a = mesh.positions[mesh.indices[t * 3]];
 			const auto& b = mesh.positions[mesh.indices[t * 3 + 1]];
@@ -488,9 +595,8 @@ namespace UgcHsr {
 			// A path from each point, then another from each, ...: a triangle that's seen is usually known at once
 			for (int sample = 0; sample < samples && !escaped; sample++) {
 				for (size_t i = 0; i < weights.size(); i++) {
-					Random random(PathSeed(options.seed, t, i, static_cast<uint64_t>(sample)));
 					paths++;
-					if (tracer.Escapes(t, weights[i], random)) {
+					if (tracer.Escapes(t, weights[i], PathSeed(options.seed, t, i, static_cast<uint64_t>(sample)))) {
 						escaped = true;
 						break;
 					}
--- a/dUgcServer/Render/UgcHsr.h
+++ b/dUgcServer/Render/UgcHsr.h
@@ -39,6 +39,7 @@ namespace UgcHsr {
 		uint64_t seed{};                // of the paths' random numbers (the same seed gives the same result)
 		UgcRays::eBackend rays{};       // what traces the paths' rays (ray_backend)
 		int fastResolution{ 1024 };     // hsr_fast_resolution: pixels square of each of the fast method's renders
+		bool sideBySide{};              // trace the paths side by side as for a GPU whatever the backend (to compare the two)
 	};

 	struct Result {
--- a/dUgcServer/Render/UgcRays.cpp
+++ b/dUgcServer/Render/UgcRays.cpp
@@ -12,6 +12,10 @@

 #include <embree4/rtcore.h>

+#ifdef DLU_HIPRT
+#include "UgcRaysHiprt.h"
+#endif
+
 namespace {
 	using UgcRays::Hit;
 	using UgcRays::INF;
@@ -528,6 +532,14 @@ namespace {
 }

 namespace UgcRays {
+	void Scene::Closest(const Ray* rays, Hit* hits, size_t count) const {
+		for (size_t i = 0; i < count; i++) hits[i] = Closest(rays[i].origin, rays[i].direction, rays[i].skip, rays[i].maxT);
+	}
+
+	void Scene::Occluded(const Ray* rays, uint8_t* occluded, size_t count) const {
+		for (size_t i = 0; i < count; i++) occluded[i] = Occluded(rays[i].origin, rays[i].direction, rays[i].minT, rays[i].maxT) ? 1 : 0;
+	}
+
 	std::string_view Name(eBackend backend) {
 		switch (backend) {
 		case eBackend::EMBREE: return "embree";
@@ -544,6 +556,9 @@ namespace UgcRays {
 	}

 	bool Available(eBackend backend) {
+#ifdef DLU_HIPRT
+		if (backend == eBackend::HIPRT) return UgcRaysHiprt::Available();
+#endif
 		return backend == eBackend::BUILTIN || backend == eBackend::EMBREE;
 	}

@@ -551,10 +566,32 @@ namespace UgcRays {
 		return Available(wanted) ? wanted : eBackend::EMBREE;
 	}

+	std::string Problem(eBackend backend) {
+		if (Available(backend)) return {};
+#ifdef DLU_HIPRT
+		if (backend == eBackend::HIPRT) return UgcRaysHiprt::Problem();
+#endif
+		return "the server was built without it (DLU_HIPRT)";
+	}
+
 	std::unique_ptr<Scene> Make(eBackend backend, const UgcModel::Mesh& mesh) {
 		switch (Resolve(backend)) {
+#ifdef DLU_HIPRT
+		case eBackend::HIPRT:
+			// A GPU that fails now (out of memory, ...) leaves the job to Embree
+			if (auto scene = UgcRaysHiprt::Make(mesh)) return scene;
+			return std::make_unique<EmbreeScene>(mesh);
+#endif
 		case eBackend::EMBREE: return std::make_unique<EmbreeScene>(mesh);
 		default: return std::make_unique<BuiltinScene>(mesh);
 		}
 	}
+
+	void SetGpuDevice(int index) {
+#ifdef DLU_HIPRT
+		UgcRaysHiprt::SetDevice(index);
+#else
+		(void)index;
+#endif
+	}
 }
--- a/dUgcServer/Render/UgcRays.h
+++ b/dUgcServer/Render/UgcRays.h
@@ -4,6 +4,7 @@
 #include <limits>
 #include <memory>
 #include <optional>
+#include <string>
 #include <string_view>

 #include <glm/glm.hpp>
@@ -15,6 +16,9 @@
 * (whether anything is hit), by one of several backends (the ugc_ray_backend setting, or per job):
 *   builtin: the UGC server's own bounding volume hierarchies (the ones it always had)
 *   embree:  Intel's Embree 4 on the CPU, on the thread that asks (no threads of its own)
+ *   hiprt:   AMD's HIPRT on the GPU (AMD through HIP, NVIDIA through CUDA, loaded when first asked for by Orochi),
+ *            when built with DLU_HIPRT and a GPU is there; else embree. One GPU for the process, used by one thread
+ *            at a time; its time is not CPU time.
 * A scene is built and traced on the thread that asks, so its time counts towards that thread's CPU time
 * (UgcThrottle). Scenes aren't shared between threads. docs/UgcServer.md ("Processing options") has the details.
 */
@@ -32,6 +36,8 @@ namespace UgcRays {
 	bool Available(eBackend backend);
 	// The backend that is used when `wanted` is asked for: itself, or embree when it isn't available
 	eBackend Resolve(eBackend wanted);
+	// Why a backend isn't available (empty when it is)
+	std::string Problem(eBackend backend);

 	struct Hit {
 		float t{ INF };
@@ -39,6 +45,18 @@ namespace UgcRays {
 		float u{}, v{}; // weights of the triangle's second and third vertex
 	};

+	// A ray of a batch (the layout the GPU kernels read too)
+	struct Ray {
+		glm::vec3 origin{};
+		float minT{};            // Occluded: hits further than this count (Closest: further than 0)
+		glm::vec3 direction{};   // unit
+		float maxT{ INF };       // hits nearer than this count
+		uint32_t skip{ NONE };   // Closest: the triangle never hit (the one the ray leaves)
+		uint32_t padding[3]{};
+	};
+	static_assert(sizeof(Ray) == 48, "the GPU kernels read rays as 48 bytes");
+	static_assert(sizeof(Hit) == 16, "the GPU kernels write hits as 16 bytes");
+
 	class Scene {
 	public:
 		virtual ~Scene() = default;
@@ -49,6 +67,13 @@ namespace UgcRays {

 		// Whether the ray (unit direction) hits a triangle further than `minT` and nearer than `maxT`
 		virtual bool Occluded(const glm::vec3& origin, const glm::vec3& direction, float minT, float maxT) const = 0;
+
+		// Many rays at once, as the single ray queries answer them (a GPU answers a batch at the cost of one ray)
+		virtual void Closest(const Ray* rays, Hit* hits, size_t count) const;
+		virtual void Occluded(const Ray* rays, uint8_t* occluded, size_t count) const;
+
+		// Whether the backend is only fast with big batches (a GPU): the callers then trace many paths side by side
+		virtual bool PrefersBatches() const { return false; }
 	};

 	/**
@@ -56,4 +81,7 @@ namespace UgcRays {
 	 * when the backend fails.
 	 */
 	std::unique_ptr<Scene> Make(eBackend backend, const UgcModel::Mesh& mesh);
+
+	// Which GPU hiprt uses (hiprt_device: 0 is the first HIP or CUDA device); before it is first used
+	void SetGpuDevice(int index);
 }
--- a/dUgcServer/Render/UgcRaysHiprt.cpp
+++ b/dUgcServer/Render/UgcRaysHiprt.cpp
@@ -0,0 +1,310 @@
+#include "UgcRaysHiprt.h"
+
+#include <algorithm>
+#include <atomic>
+#include <filesystem>
+#include <mutex>
+#include <stdexcept>
+#include <vector>
+
+#include <Orochi/Orochi.h>
+
+// hiprtew.h loads HIPRT's library when asked (hiprtewInit) instead of linking it; its function pointers live here
+#ifndef _WIN32
+#define HIPRTAPI
+#endif
+#define _ENABLE_HIPRTEW
+#include <hiprt/hiprtew.h>
+
+#include "BinaryPathFinder.h"
+
+namespace {
+	// The trace kernels, compiled at run time by HIPRT (for the GPU found). A ray is 48 bytes and a hit 16, as
+	// UgcRays::Ray and UgcRays::Hit. Closest never returns the ray's `skip` triangle: when it is the nearest, the
+	// ray goes on from just past it (a few times at most).
+	const char* KERNELS = R"(
+#include <hiprt/hiprt_device.h>
+
+struct UgcRay { float ox, oy, oz, minT, dx, dy, dz, maxT; unsigned skip, pad0, pad1, pad2; };
+struct UgcHit { float t; unsigned triangle; float u, v; };
+
+extern "C" __global__ void UgcClosest(hiprtGeometry geometry, const UgcRay* rays, UgcHit* hits, unsigned count) {
+	const unsigned i = blockIdx.x * blockDim.x + threadIdx.x;
+	if (i >= count) return;
+	const UgcRay r = rays[i];
+	hiprtRay ray;
+	ray.origin = { r.ox, r.oy, r.oz };
+	ray.direction = { r.dx, r.dy, r.dz };
+	ray.minT = 1.17549435e-38f;
+	// Finite: a ray grazing a triangle can otherwise hit it at infinity
+	ray.maxT = fminf(r.maxT, 3.40282347e38f);
+	UgcHit out = { r.maxT, 0xFFFFFFFFu, 0.0f, 0.0f };
+	for (int attempt = 0; attempt < 8; attempt++) {
+		hiprtGeomTraversalClosest traversal(geometry, ray);
+		const hiprtHit hit = traversal.getNextHit();
+		if (!hit.hasHit()) break;
+		if (hit.primID != r.skip) {
+			out.t = hit.t;
+			out.triangle = hit.primID;
+			out.u = hit.uv.x;
+			out.v = hit.uv.y;
+			break;
+		}
+		ray.minT = nextafterf(hit.t, 3.4e38f);
+	}
+	hits[i] = out;
+}
+
+extern "C" __global__ void UgcOccluded(hiprtGeometry geometry, const UgcRay* rays, unsigned char* occluded, unsigned count) {
+	const unsigned i = blockIdx.x * blockDim.x + threadIdx.x;
+	if (i >= count) return;
+	const UgcRay r = rays[i];
+	hiprtRay ray;
+	ray.origin = { r.ox, r.oy, r.oz };
+	ray.direction = { r.dx, r.dy, r.dz };
+	// Hits exactly at minT or maxT don't count, as in the other backends
+	ray.minT = nextafterf(r.minT, 3.4e38f);
+	ray.maxT = nextafterf(fminf(r.maxT, 3.40282347e38f), 0.0f);
+	hiprtGeomTraversalAnyHit traversal(geometry, ray);
+	occluded[i] = traversal.getNextHit().hasHit() ? 1 : 0;
+}
+)";
+
+	// Rays sent to the GPU at once at most (48 MB of rays)
+	constexpr size_t BATCH = 1u << 20;
+	constexpr unsigned BLOCK = 64;
+
+	std::atomic<int> g_DeviceIndex{ 0 };
+	std::mutex g_Mutex; // one thread on the GPU at a time; everything below is used under it
+
+	struct Gpu {
+		bool tried{};
+		bool ok{};
+		std::string problem;
+		oroDevice device{};
+		oroCtx context{};
+		hiprtContext rt{};
+		hiprtApiFunction closest{};
+		hiprtApiFunction occluded{};
+		void* rays{};   // device buffers for a batch, grown as needed
+		void* results{};
+		size_t capacity{};
+	} g_Gpu;
+
+	std::string OroMessage(oroError error) {
+		const char* text = nullptr;
+		oroGetErrorString(error, &text);
+		return text ? text : "error " + std::to_string(static_cast<int>(error));
+	}
+
+	void Check(oroError error, const char* what) {
+		if (error != oroSuccess) throw std::runtime_error(std::string("GPU: ") + what + " failed (" + OroMessage(error) + ")");
+	}
+
+	void Check(hiprtError error, const char* what) {
+		if (error != hiprtSuccess) throw std::runtime_error(std::string("HIPRT: ") + what + " failed (error " + std::to_string(static_cast<int>(error)) + ")");
+	}
+
+	// Where the HIPRT headers the kernels include are: next to the servers (copied there by the build), else where
+	// the build found them
+	std::string IncludeFolder() {
+		const auto local = BinaryPathFinder::GetBinaryDir() / "hiprt" / "include";
+		if (std::filesystem::exists(local / "hiprt" / "hiprt_device.h")) return local.string();
+		return DLU_HIPRT_INCLUDE_DIR;
+	}
+
+	// Under g_Mutex: loads HIP or CUDA and HIPRT, makes the context and compiles the kernels, once
+	bool Init() {
+		if (g_Gpu.tried) return g_Gpu.ok;
+		g_Gpu.tried = true;
+		try {
+			if (oroInitialize(static_cast<oroApi>(ORO_API_HIP | ORO_API_CUDA), 0) != 0) throw std::runtime_error("neither HIP nor CUDA could be loaded");
+			Check(oroInit(0), "oroInit");
+			int count = 0;
+			Check(oroGetDeviceCount(&count), "oroGetDeviceCount");
+			const int index = g_DeviceIndex;
+			if (index < 0 || index >= count) throw std::runtime_error("no GPU " + std::to_string(index) + " (" + std::to_string(count) + " found)");
+			Check(oroDeviceGet(&g_Gpu.device, index), "oroDeviceGet");
+			Check(oroCtxCreate(&g_Gpu.context, 0, g_Gpu.device), "oroCtxCreate");
+			int loaded = 0;
+			hiprtewInit(&loaded);
+			if (loaded != HIPRTEW_SUCCESS) throw std::runtime_error(std::string("the HIPRT library (") + HIPRT_LIB_NAME + ") could not be loaded");
+			hiprtContextCreationInput input{};
+			input.ctxt = oroGetRawCtx(g_Gpu.context);
+			input.device = oroGetRawDevice(g_Gpu.device);
+			input.deviceType = oroGetCurAPI(0) == ORO_API_CUDADRIVER ? hiprtDeviceNVIDIA : hiprtDeviceAMD;
+			Check(hiprtCreateContext(HIPRT_API_VERSION, input, g_Gpu.rt), "hiprtCreateContext");
+			const auto cache = BinaryPathFinder::GetBinaryDir() / "cache" / "hiprt";
+			std::error_code error;
+			std::filesystem::create_directories(cache, error);
+			if (!error) hiprtSetCacheDirPath(g_Gpu.rt, cache.string().c_str());
+			const char* names[] = { "UgcClosest", "UgcOccluded" };
+			hiprtApiFunction functions[2]{};
+			const auto include = "-I" + IncludeFolder();
+			const char* options[] = { include.c_str() };
+			Check(hiprtBuildTraceKernels(g_Gpu.rt, 2, names, KERNELS, "UgcRays", 0, nullptr, nullptr, 1, options, 0, 1, nullptr, functions, nullptr, true),
+				"hiprtBuildTraceKernels");
+			g_Gpu.closest = functions[0];
+			g_Gpu.occluded = functions[1];
+			g_Gpu.ok = true;
+		} catch (const std::exception& ex) {
+			g_Gpu.problem = ex.what();
+			g_Gpu.ok = false;
+		}
+		return g_Gpu.ok;
+	}
+
+	// Under g_Mutex, with the context current: device buffers for `count` rays and their results
+	void Reserve(size_t count) {
+		if (count <= g_Gpu.capacity) return;
+		if (g_Gpu.rays) oroFree(g_Gpu.rays);
+		if (g_Gpu.results) oroFree(g_Gpu.results);
+		g_Gpu.rays = g_Gpu.results = nullptr;
+		g_Gpu.capacity = 0;
+		Check(oroMalloc(&g_Gpu.rays, count * sizeof(UgcRays::Ray)), "oroMalloc");
+		Check(oroMalloc(&g_Gpu.results, count * sizeof(UgcRays::Hit)), "oroMalloc");
+		g_Gpu.capacity = count;
+	}
+
+	class HiprtScene final : public UgcRays::Scene {
+	public:
+		explicit HiprtScene(const UgcModel::Mesh& mesh) {
+			m_Triangles = static_cast<uint32_t>(mesh.TriangleCount());
+			if (m_Triangles == 0 || mesh.positions.empty()) return;
+			std::lock_guard lock(g_Mutex);
+			Check(oroCtxSetCurrent(g_Gpu.context), "oroCtxSetCurrent");
+			try {
+				const size_t vertexBytes = mesh.positions.size() * sizeof(glm::vec3), indexBytes = static_cast<size_t>(m_Triangles) * 3 * sizeof(uint32_t);
+				Check(oroMalloc(&m_Vertices, vertexBytes), "oroMalloc");
+				Check(oroMalloc(&m_Indices, indexBytes), "oroMalloc");
+				Check(oroMemcpyHtoD(reinterpret_cast<oroDeviceptr>(m_Vertices), const_cast<glm::vec3*>(mesh.positions.data()), vertexBytes), "oroMemcpyHtoD");
+				Check(oroMemcpyHtoD(reinterpret_cast<oroDeviceptr>(m_Indices), const_cast<uint32_t*>(mesh.indices.data()), indexBytes), "oroMemcpyHtoD");
+				hiprtTriangleMeshPrimitive triangles{};
+				triangles.vertices = m_Vertices;
+				triangles.vertexCount = static_cast<uint32_t>(mesh.positions.size());
+				triangles.vertexStride = sizeof(glm::vec3);
+				triangles.triangleIndices = m_Indices;
+				triangles.triangleCount = m_Triangles;
+				triangles.triangleStride = 3 * sizeof(uint32_t);
+				hiprtGeometryBuildInput input{};
+				input.type = hiprtPrimitiveTypeTriangleMesh;
+				input.primitive.triangleMesh = triangles;
+				hiprtBuildOptions options{};
+				options.buildFlags = hiprtBuildFlagBitPreferHighQualityBuild;
+				size_t temporaryBytes = 0;
+				Check(hiprtGetGeometryBuildTemporaryBufferSize(g_Gpu.rt, input, options, temporaryBytes), "hiprtGetGeometryBuildTemporaryBufferSize");
+				void* temporary = nullptr;
+				if (temporaryBytes > 0) Check(oroMalloc(&temporary, temporaryBytes), "oroMalloc");
+				Check(hiprtCreateGeometry(g_Gpu.rt, input, options, m_Geometry), "hiprtCreateGeometry");
+				m_Created = true;
+				const auto built = hiprtBuildGeometry(g_Gpu.rt, hiprtBuildOperationBuild, input, options, temporary, nullptr, m_Geometry);
+				if (temporary) oroFree(temporary);
+				Check(built, "hiprtBuildGeometry");
+			} catch (...) {
+				Release();
+				throw;
+			}
+		}
+
+		~HiprtScene() override {
+			std::lock_guard lock(g_Mutex);
+			if (oroCtxSetCurrent(g_Gpu.context) == oroSuccess) Release();
+		}
+
+		HiprtScene(const HiprtScene&) = delete;
+		HiprtScene& operator=(const HiprtScene&) = delete;
+
+		UgcRays::Hit Closest(const glm::vec3& origin, const glm::vec3& direction, uint32_t skip, float maxT) const override {
+			UgcRays::Ray ray{ origin, 0.0f, direction, maxT, skip };
+			UgcRays::Hit hit;
+			Closest(&ray, &hit, 1);
+			return hit;
+		}
+
+		bool Occluded(const glm::vec3& origin, const glm::vec3& direction, float minT, float maxT) const override {
+			UgcRays::Ray ray{ origin, minT, direction, maxT };
+			uint8_t occluded = 0;
+			Occluded(&ray, &occluded, 1);
+			return occluded != 0;
+		}
+
+		void Closest(const UgcRays::Ray* rays, UgcRays::Hit* hits, size_t count) const override {
+			if (m_Triangles == 0) {
+				for (size_t i = 0; i < count; i++) hits[i] = UgcRays::Hit{ rays[i].maxT };
+				return;
+			}
+			Trace(g_Gpu.closest, rays, hits, sizeof(UgcRays::Hit), count);
+		}
+
+		void Occluded(const UgcRays::Ray* rays, uint8_t* occluded, size_t count) const override {
+			if (m_Triangles == 0) {
+				std::fill(occluded, occluded + count, uint8_t{ 0 });
+				return;
+			}
+			Trace(g_Gpu.occluded, rays, occluded, sizeof(uint8_t), count);
+		}
+
+		bool PrefersBatches() const override { return true; }
+
+	private:
+		// The rays through `function` in batches; `results` gets `resultBytes` a ray
+		void Trace(hiprtApiFunction function, const UgcRays::Ray* rays, void* results, size_t resultBytes, size_t count) const {
+			std::lock_guard lock(g_Mutex);
+			Check(oroCtxSetCurrent(g_Gpu.context), "oroCtxSetCurrent");
+			Reserve(std::min(count, BATCH));
+			for (size_t first = 0; first < count; first += BATCH) {
+				const auto batch = static_cast<unsigned>(std::min(BATCH, count - first));
+				Check(oroMemcpyHtoD(reinterpret_cast<oroDeviceptr>(g_Gpu.rays), const_cast<UgcRays::Ray*>(rays + first), batch * sizeof(UgcRays::Ray)), "oroMemcpyHtoD");
+				hiprtGeometry geometry = m_Geometry;
+				void* deviceRays = g_Gpu.rays;
+				void* deviceResults = g_Gpu.results;
+				unsigned batchCount = batch;
+				void* arguments[] = { &geometry, &deviceRays, &deviceResults, &batchCount };
+				Check(oroModuleLaunchKernel(reinterpret_cast<oroFunction>(function), (batch + BLOCK - 1) / BLOCK, 1, 1, BLOCK, 1, 1, 0, nullptr, arguments, nullptr),
+					"oroModuleLaunchKernel");
+				Check(oroDeviceSynchronize(), "oroDeviceSynchronize");
+				Check(oroMemcpyDtoH(static_cast<uint8_t*>(results) + first * resultBytes, reinterpret_cast<oroDeviceptr>(g_Gpu.results), batch * resultBytes), "oroMemcpyDtoH");
+			}
+		}
+
+		// Under g_Mutex with the context current
+		void Release() {
+			if (m_Created) hiprtDestroyGeometry(g_Gpu.rt, m_Geometry);
+			if (m_Vertices) oroFree(m_Vertices);
+			if (m_Indices) oroFree(m_Indices);
+			m_Created = false;
+			m_Vertices = m_Indices = nullptr;
+		}
+
+		uint32_t m_Triangles{};
+		void* m_Vertices{};
+		void* m_Indices{};
+		hiprtGeometry m_Geometry{};
+		bool m_Created{};
+	};
+}
+
+namespace UgcRaysHiprt {
+	bool Available() {
+		std::lock_guard lock(g_Mutex);
+		return Init();
+	}
+
+	std::string Problem() {
+		std::lock_guard lock(g_Mutex);
+		return g_Gpu.problem;
+	}
+
+	std::unique_ptr<UgcRays::Scene> Make(const UgcModel::Mesh& mesh) {
+		if (!Available()) return nullptr;
+		try {
+			return std::make_unique<HiprtScene>(mesh);
+		} catch (const std::exception&) {
+			return nullptr;
+		}
+	}
+
+	void SetDevice(int index) {
+		g_DeviceIndex = index;
+	}
+}
--- a/dUgcServer/Render/UgcRaysHiprt.h
+++ b/dUgcServer/Render/UgcRaysHiprt.h
@@ -0,0 +1,25 @@
+#pragma once
+
+#include <memory>
+#include <string>
+
+#include "UgcRays.h"
+
+/**
+ * UgcRays' hiprt backend (built with DLU_HIPRT): AMD's HIPRT on the GPU, with HIP or CUDA loaded at run time through
+ * Orochi. One GPU context for the process; the threads take turns on it. The trace kernels are compiled the first time
+ * (HIPRT keeps them in its cache folder after that).
+ */
+namespace UgcRaysHiprt {
+	// Whether the GPU can be used: loads HIP or CUDA and HIPRT, makes the context and compiles the kernels the first time
+	bool Available();
+
+	// Why it can't (empty when it can, or before it was first asked)
+	std::string Problem();
+
+	// The mesh on the GPU; null when the GPU fails (the caller uses another backend)
+	std::unique_ptr<UgcRays::Scene> Make(const UgcModel::Mesh& mesh);
+
+	// Which device (0: the first); before it is first used
+	void SetDevice(int index);
+}
--- a/dUgcServer/Render/UgcRender.cpp
+++ b/dUgcServer/Render/UgcRender.cpp
@@ -284,36 +284,60 @@ namespace UgcRender {
 				return h;
 			}
 		};
-		std::unordered_map<Key, float, KeyHash> known;
-		known.reserve(mesh.positions.size());
+		// Which vertices are worked out (the first at each place and facing) and which take another's
+		std::unordered_map<Key, size_t, KeyHash> first;
+		first.reserve(mesh.positions.size());
+		std::vector<size_t> copyOf(mesh.positions.size(), SIZE_MAX);
+		std::vector<size_t> traced;
 		for (size_t v = 0; v < mesh.positions.size(); v++) {
-			if ((v & 0xFF) == 0) UgcThrottle::Checkpoint();
 			const auto& normal = mesh.normals[v];
 			const Key key{ { static_cast<int32_t>(std::lround(mesh.positions[v].x * 1000.0f)), static_cast<int32_t>(std::lround(mesh.positions[v].y * 1000.0f)),
 				static_cast<int32_t>(std::lround(mesh.positions[v].z * 1000.0f)) }, { static_cast<int32_t>(std::lround(normal.x * 100.0f)),
 				static_cast<int32_t>(std::lround(normal.y * 100.0f)), static_cast<int32_t>(std::lround(normal.z * 100.0f)) } };
-			if (const auto it = known.find(key); it != known.end()) {
-				ao[v] = it->second;
+			if (const auto it = first.find(key); it != first.end()) {
+				copyOf[v] = it->second;
 				continue;
 			}
 			if (glm::dot(normal, normal) < 0.5f) continue;
-			// A frame around the normal
-			const glm::vec3 helper = std::abs(normal.x) < 0.9f ? glm::vec3(1, 0, 0) : glm::vec3(0, 1, 0);
-			const auto tangent = glm::normalize(glm::cross(helper, normal));
-			const auto bitangent = glm::cross(normal, tangent);
-			// Hammersley points, turned by an amount of the vertex's own (fixed) so neighbours don't band
-			const float turn = static_cast<float>((v * 0x9E3779B9u) >> 8 & 0xFFFFFF) / 16777216.0f;
-			const auto origin = mesh.positions[v] + normal * 1e-3f;
-			uint32_t open = 0;
-			for (uint32_t i = 0; i < count; i++) {
-				const float u = (i + 0.5f) / static_cast<float>(count);
-				const float phi = 2.0f * 3.14159265f * std::fmod(RadicalInverse(i) + turn, 1.0f);
-				const float r = std::sqrt(u), z = std::sqrt(std::max(0.0f, 1.0f - u));
-				const auto direction = tangent * (r * std::cos(phi)) + bitangent * (r * std::sin(phi)) + normal * z;
-				if (!scene->Occluded(origin, direction, 1e-4f, distance)) open++;
+			first.emplace(key, v);
+			traced.push_back(v);
+		}
+		// The rays in batches of vertices: small ones between checkpoints on the CPU, big ones for a GPU
+		const size_t perBatch = scene->PrefersBatches() ? std::max<size_t>(1, (1u << 18) / count) : 256;
+		std::vector<UgcRays::Ray> batch;
+		std::vector<uint8_t> occluded;
+		for (size_t start = 0; start < traced.size(); start += perBatch) {
+			UgcThrottle::Checkpoint();
+			const size_t end = std::min(traced.size(), start + perBatch);
+			batch.clear();
+			for (size_t j = start; j < end; j++) {
+				const auto v = traced[j];
+				const auto& normal = mesh.normals[v];
+				// A frame around the normal
+				const glm::vec3 helper = std::abs(normal.x) < 0.9f ? glm::vec3(1, 0, 0) : glm::vec3(0, 1, 0);
+				const auto tangent = glm::normalize(glm::cross(helper, normal));
+				const auto bitangent = glm::cross(normal, tangent);
+				// Hammersley points, turned by an amount of the vertex's own (fixed) so neighbours don't band
+				const float turn = static_cast<float>((v * 0x9E3779B9u) >> 8 & 0xFFFFFF) / 16777216.0f;
+				const auto origin = mesh.positions[v] + normal * 1e-3f;
+				for (uint32_t i = 0; i < count; i++) {
+					const float u = (i + 0.5f) / static_cast<float>(count);
+					const float phi = 2.0f * 3.14159265f * std::fmod(RadicalInverse(i) + turn, 1.0f);
+					const float r = std::sqrt(u), z = std::sqrt(std::max(0.0f, 1.0f - u));
+					const auto direction = tangent * (r * std::cos(phi)) + bitangent * (r * std::sin(phi)) + normal * z;
+					batch.push_back(UgcRays::Ray{ origin, 1e-4f, direction, distance });
+				}
 			}
-			ao[v] = static_cast<float>(open) / static_cast<float>(count);
-			known.emplace(key, ao[v]);
+			occluded.resize(batch.size());
+			scene->Occluded(batch.data(), occluded.data(), batch.size());
+			for (size_t j = start; j < end; j++) {
+				uint32_t open = 0;
+				for (uint32_t i = 0; i < count; i++) open += occluded[(j - start) * count + i] ? 0 : 1;
+				ao[traced[j]] = static_cast<float>(open) / static_cast<float>(count);
+			}
+		}
+		for (size_t v = 0; v < ao.size(); v++) {
+			if (copyOf[v] != SIZE_MAX) ao[v] = ao[copyOf[v]];
 		}
 		return ao;
 	}
@@ -517,31 +541,49 @@ namespace UgcRender {
 		if (traced) {
 			const auto scene = UgcRays::Make(options.ao.rays, model.opaque);
 			const auto count = static_cast<uint32_t>(std::clamp(options.denoiseSamples, 1, 256));
+			// In batches of pixels: small ones between checkpoints on the CPU, big ones for a GPU
+			const size_t perBatch = scene->PrefersBatches() ? std::max<size_t>(1, (1u << 18) / count) : 1024;
+			std::vector<size_t> drawn;
 			for (size_t index = 0; index < color.size(); index++) {
-				if ((index & 0x3FF) == 0) UgcThrottle::Checkpoint();
-				if (color[index].a <= 0.0f) continue;
-				const auto& normal = pointNormals[index];
-				const glm::vec3 helper = std::abs(normal.x) < 0.9f ? glm::vec3(1, 0, 0) : glm::vec3(0, 1, 0);
-				const auto tangent = glm::normalize(glm::cross(helper, normal));
-				const auto bitangent = glm::cross(normal, tangent);
-				uint32_t hash = static_cast<uint32_t>(index) * 0x9E3779B9u;
-				hash ^= hash >> 16;
-				hash *= 0x85EBCA6Bu;
-				hash ^= hash >> 13;
-				const float turn = static_cast<float>(hash >> 8) / 16777216.0f;
-				const float shift = static_cast<float>((hash * 0xC2B2AE35u) >> 8) / 16777216.0f;
-				const auto origin = points[index] + normal * 1e-3f;
-				uint32_t open = 0;
-				for (uint32_t i = 0; i < count; i++) {
-					const float u = std::fmod((i + shift) / static_cast<float>(count), 1.0f);
-					const float phi = 2.0f * 3.14159265f * std::fmod(RadicalInverse(i) + turn, 1.0f);
-					const float r = std::sqrt(u), up = std::sqrt(std::max(0.0f, 1.0f - u));
-					const auto direction = tangent * (r * std::cos(phi)) + bitangent * (r * std::sin(phi)) + normal * up;
-					if (!scene->Occluded(origin, direction, 1e-4f, options.ao.distance)) open++;
+				if (color[index].a > 0.0f) drawn.push_back(index);
+			}
+			std::vector<UgcRays::Ray> batch;
+			std::vector<uint8_t> occluded;
+			for (size_t start = 0; start < drawn.size(); start += perBatch) {
+				UgcThrottle::Checkpoint();
+				const size_t end = std::min(drawn.size(), start + perBatch);
+				batch.clear();
+				for (size_t j = start; j < end; j++) {
+					const auto index = drawn[j];
+					const auto& normal = pointNormals[index];
+					const glm::vec3 helper = std::abs(normal.x) < 0.9f ? glm::vec3(1, 0, 0) : glm::vec3(0, 1, 0);
+					const auto tangent = glm::normalize(glm::cross(helper, normal));
+					const auto bitangent = glm::cross(normal, tangent);
+					uint32_t hash = static_cast<uint32_t>(index) * 0x9E3779B9u;
+					hash ^= hash >> 16;
+					hash *= 0x85EBCA6Bu;
+					hash ^= hash >> 13;
+					const float turn = static_cast<float>(hash >> 8) / 16777216.0f;
+					const float shift = static_cast<float>((hash * 0xC2B2AE35u) >> 8) / 16777216.0f;
+					const auto origin = points[index] + normal * 1e-3f;
+					for (uint32_t i = 0; i < count; i++) {
+						const float u = std::fmod((i + shift) / static_cast<float>(count), 1.0f);
+						const float phi = 2.0f * 3.14159265f * std::fmod(RadicalInverse(i) + turn, 1.0f);
+						const float r = std::sqrt(u), up = std::sqrt(std::max(0.0f, 1.0f - u));
+						const auto direction = tangent * (r * std::cos(phi)) + bitangent * (r * std::sin(phi)) + normal * up;
+						batch.push_back(UgcRays::Ray{ origin, 1e-4f, direction, options.ao.distance });
+					}
+				}
+				occluded.resize(batch.size());
+				scene->Occluded(batch.data(), occluded.data(), batch.size());
+				for (size_t j = start; j < end; j++) {
+					uint32_t open = 0;
+					for (uint32_t i = 0; i < count; i++) open += occluded[(j - start) * count + i] ? 0 : 1;
+					const auto index = drawn[j];
+					const float occlusion = static_cast<float>(open) / static_cast<float>(count);
+					const float lit = 1.0f - std::clamp(options.bakedAo, 0.0f, 1.0f) * (1.0f - occlusion);
+					color[index] = glm::vec4(glm::vec3(color[index]) * lit, color[index].a);
 				}
-				const float occlusion = static_cast<float>(open) / static_cast<float>(count);
-				const float lit = 1.0f - std::clamp(options.bakedAo, 0.0f, 1.0f) * (1.0f - occlusion);
-				color[index] = glm::vec4(glm::vec3(color[index]) * lit, color[index].a);
 			}
 		}

--- a/dUgcServer/UgcServer.cpp
+++ b/dUgcServer/UgcServer.cpp
@@ -164,6 +164,8 @@ namespace {
 		// What traces the rays of the hidden faces' paths and of the occlusion (the icon's too)
 		settings.hsr.rays = UgcRays::Parse(Game::config->GetValue("ray_backend")).value_or(UgcRays::eBackend::BUILTIN);
 		settings.ao.rays = settings.hsr.rays;
+		// The GPU hiprt uses (read before it is first used; changing it takes a restart)
+		UgcRays::SetGpuDevice(std::max(Setting<int32_t>("hiprt_device", 0), 0));
 		settings.ao.enabled = Setting<int32_t>("bake_ao", 1) != 0;
 		settings.ao.distance = Setting<float>("ao_distance", 5.0f);
 		settings.ao.samples = std::clamp(Setting<int32_t>("ao_samples", 64), 1, 1024);
@@ -186,6 +188,17 @@ namespace {
 		return settings;
 	}

+	// The processing options the settings pick, and what is used instead when this build or machine can't
+	// (main thread: the GPU is set up here the first time)
+	void LogProcessingOptions(const UgcJobs::Settings& settings) {
+		const auto made = UgcJobs::MadeWith(settings);
+		LOG("Processing options: %s", UgcProcessOptions::ToString(made).c_str());
+		if (UgcRays::Resolve(settings.hsr.rays) != settings.hsr.rays) {
+			LOG("ray_backend=%s can't be used (%s): embree instead", std::string(UgcRays::Name(settings.hsr.rays)).c_str(), UgcRays::Problem(settings.hsr.rays).c_str());
+		}
+		if (!UgcRender::Available(settings.icon.denoise)) LOG("denoise=%s can't be used (the server was built without DLU_OIDN): off instead", std::string(UgcRender::Name(settings.icon.denoise)).c_str());
+	}
+
 	UgcProcessor::Limits ReadLimits() {
 		UgcProcessor::Limits limits;
 		const auto cores = std::max<unsigned>(std::thread::hardware_concurrency(), 1);
@@ -691,6 +704,7 @@ int main(int argc, char** argv) {
 	processorConfig.threads = threads > 0 ? threads : std::max<size_t>(std::thread::hardware_concurrency() / 2, 1);
 	UgcProcessor processor(processorConfig, storage, library, ReadSettings());
 	processor.Configure(ReadSettings(), ReadLimits());
+	LogProcessingOptions(ReadSettings());
 	g_Processor = &processor;
 	// Sent with the traffic reports to the dashboard (Diagnostics)
 	TrafficStats::Local().SetGauge("workers_busy", [&processor] { return static_cast<double>(processor.Busy()); });
--- a/resources/ugcconfig.ini
+++ b/resources/ugcconfig.ini
@@ -146,8 +146,10 @@ hsr_method=toolbox
 hsr_fast_resolution=1024

 # What traces the rays of the hidden faces' paths and of the occlusion: builtin, embree (Intel Embree on the CPU) or
-# hiprt (the GPU, when built with DLU_HIPRT and the machine has one; else embree).
+# hiprt (the GPU, when built with DLU_HIPRT and the machine has one; else embree). hiprt_device: which GPU (0: the first
+# HIP or CUDA device; restart to change).
 ray_backend=builtin
+hiprt_device=0

 # denoise: off, or oidn (when built with DLU_OIDN; else off): a model's icon is drawn from its colors before the
 # occlusion bake with its occlusion traced per pixel (denoise_samples rays from each pixel of the supersampled image)
--- a/tests/dUgcTests/UgcTests.cpp
+++ b/tests/dUgcTests/UgcTests.cpp
@@ -2215,6 +2215,24 @@ TEST(UgcRays, BackendsFindTheSameHits) {
 	}
 }

+TEST(UgcHsr, SideBySideDecidesTheSameAsOneByOne) {
+	// The paths traced side by side (as for a GPU) are the same paths, so the same triangles stay; ground plane too
+	for (const bool ground : { false, true }) {
+		UgcHsr::Options options;
+		options.seed = 5;
+		options.samples = 2;
+		options.groundPlane = ground;
+		const auto mesh = Clutter();
+		uint64_t points = 0, paths = 0, sidePoints = 0, sidePaths = 0;
+		const auto oneByOne = UgcHsr::Visible(mesh, options, &points, &paths);
+		options.sideBySide = true;
+		EXPECT_EQ(UgcHsr::Visible(mesh, options, &sidePoints, &sidePaths), oneByOne) << ground;
+		EXPECT_EQ(sidePoints, points);
+		EXPECT_GE(sidePaths, paths); // it traces a round's other paths too once one escaped
+		EXPECT_EQ(UgcHsr::Visible(Room(true), options), UgcHsr::Visible(Room(true), UgcHsr::Options{ .groundPlane = ground, .samples = 2, .seed = 5 }));
+	}
+}
+
 TEST(UgcRays, OtherBackendsMakeTheSameModels) {
 	// The hidden faces and the occlusion with each backend. The paths bounce, so a hit found a rounding further away
 	// sends a path on from a slightly different point: the rare triangle decided by a path that only just gets out
--- a/thirdparty/CMakeLists.txt
+++ b/thirdparty/CMakeLists.txt
@@ -208,3 +208,34 @@ if(DLU_OIDN)
 	endif()
 	message(STATUS "Open Image Denoise ${OpenImageDenoise_VERSION}: the UGC server can denoise icons")
 endif()
+
+# HIPRT (MIT), the UGC server's ray_backend=hiprt on the GPU: optional, off by default. Needs HIPRT's SDK (its headers;
+# HIPRT_ROOT, else ROCm's /opt/rocm), whose library is loaded at run time (hiprtew), as HIP or CUDA are by Orochi (MIT,
+# fetched here; CUDA too when its toolkit is found). The headers are copied next to the servers: the GPU kernels are
+# compiled from them when the UGC server first uses the GPU.
+option(DLU_HIPRT "Build the UGC server with HIPRT (the ray_backend=hiprt setting, on the GPU)" OFF)
+if(DLU_HIPRT)
+	find_path(DLU_HIPRT_INCLUDE_DIR hiprt/hiprtew.h PATHS "$ENV{HIPRT_ROOT}/include" "${HIPRT_ROOT}/include" "/opt/rocm/include" REQUIRED)
+	FetchContent_Declare(orochi
+		URL https://github.com/GPUOpen-LibrariesAndSDKs/Orochi/archive/dbaf54bb61007d7ce87802b9458df54eeb574fe1.tar.gz
+		URL_HASH SHA256=13d5b946aceadc07b09633e5cb2917433209f562a596e76181325d04cd0d4173
+		DOWNLOAD_EXTRACT_TIMESTAMP TRUE)
+	FetchContent_MakeAvailable(orochi) # no CMakeLists.txt of its own: only downloaded
+	add_library(dlu_orochi STATIC "${orochi_SOURCE_DIR}/Orochi/Orochi.cpp" "${orochi_SOURCE_DIR}/contrib/hipew/src/hipew.cpp")
+	target_include_directories(dlu_orochi PUBLIC "${orochi_SOURCE_DIR}")
+	target_link_libraries(dlu_orochi PUBLIC ${CMAKE_DL_LIBS})
+	# The servers export their symbols (-rdynamic): HIPRT's library has an Orochi of its own, whose names ours mustn't take
+	set_target_properties(dlu_orochi PROPERTIES CXX_VISIBILITY_PRESET hidden VISIBILITY_INLINES_HIDDEN ON)
+	find_package(CUDAToolkit QUIET)
+	if(CUDAToolkit_FOUND)
+		target_sources(dlu_orochi PRIVATE "${orochi_SOURCE_DIR}/contrib/cuew/src/cuew.cpp")
+		target_compile_definitions(dlu_orochi PUBLIC OROCHI_ENABLE_CUEW)
+		target_include_directories(dlu_orochi PUBLIC ${CUDAToolkit_INCLUDE_DIRS})
+	endif()
+	file(COPY "${DLU_HIPRT_INCLUDE_DIR}/hiprt" DESTINATION "${CMAKE_BINARY_DIR}/hiprt/include")
+	if(CUDAToolkit_FOUND)
+		message(STATUS "HIPRT headers in ${DLU_HIPRT_INCLUDE_DIR}: the UGC server can trace rays on AMD (HIP) and NVIDIA (CUDA) GPUs")
+	else()
+		message(STATUS "HIPRT headers in ${DLU_HIPRT_INCLUDE_DIR}: the UGC server can trace rays on AMD GPUs (HIP; no CUDA toolkit found)")
+	endif()
+endif()