diff --git a/dDashboardServer/routes/SettingsCatalog.cpp b/dDashboardServer/routes/SettingsCatalog.cpp index 0ff28b4d0..9a2ae3b0a 100644 --- 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)); diff --git a/dUgcServer/CMakeLists.txt b/dUgcServer/CMakeLists.txt index d601ce58c..019f900dc 100644 --- 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") diff --git a/dUgcServer/Render/UgcHsr.cpp b/dUgcServer/Render/UgcHsr.cpp index 88373c5ad..e41d77290 100644 --- 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 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 group; + std::vector> 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(std::max(options.minPoints, 1)))); + group.push_back(next); + groupPaths += groupWeights.back().size(); + points += groupWeights.back().size(); + } + std::vector escaped(group.size(), 0); + for (int sample = 0; sample < samples; sample++) { + std::vector live; + std::vector 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(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 rays; + std::vector 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(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(sample))); paths++; - if (tracer.Escapes(t, weights[i], random)) { + if (tracer.Escapes(t, weights[i], PathSeed(options.seed, t, i, static_cast(sample)))) { escaped = true; break; } diff --git a/dUgcServer/Render/UgcHsr.h b/dUgcServer/Render/UgcHsr.h index 8c63e5763..db6fc3121 100644 --- 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 { diff --git a/dUgcServer/Render/UgcRays.cpp b/dUgcServer/Render/UgcRays.cpp index f254cfce8..7427e2f4a 100644 --- a/dUgcServer/Render/UgcRays.cpp +++ b/dUgcServer/Render/UgcRays.cpp @@ -12,6 +12,10 @@ #include +#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 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(mesh); +#endif case eBackend::EMBREE: return std::make_unique(mesh); default: return std::make_unique(mesh); } } + + void SetGpuDevice(int index) { +#ifdef DLU_HIPRT + UgcRaysHiprt::SetDevice(index); +#else + (void)index; +#endif + } } diff --git a/dUgcServer/Render/UgcRays.h b/dUgcServer/Render/UgcRays.h index 9e95f7f84..241167f58 100644 --- a/dUgcServer/Render/UgcRays.h +++ b/dUgcServer/Render/UgcRays.h @@ -4,6 +4,7 @@ #include #include #include +#include #include #include @@ -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 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); } diff --git a/dUgcServer/Render/UgcRaysHiprt.cpp b/dUgcServer/Render/UgcRaysHiprt.cpp new file mode 100644 index 000000000..bf434b71e --- /dev/null +++ b/dUgcServer/Render/UgcRaysHiprt.cpp @@ -0,0 +1,310 @@ +#include "UgcRaysHiprt.h" + +#include +#include +#include +#include +#include +#include + +#include + +// 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 + +#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 + +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 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(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(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(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(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(m_Triangles) * 3 * sizeof(uint32_t); + Check(oroMalloc(&m_Vertices, vertexBytes), "oroMalloc"); + Check(oroMalloc(&m_Indices, indexBytes), "oroMalloc"); + Check(oroMemcpyHtoD(reinterpret_cast(m_Vertices), const_cast(mesh.positions.data()), vertexBytes), "oroMemcpyHtoD"); + Check(oroMemcpyHtoD(reinterpret_cast(m_Indices), const_cast(mesh.indices.data()), indexBytes), "oroMemcpyHtoD"); + hiprtTriangleMeshPrimitive triangles{}; + triangles.vertices = m_Vertices; + triangles.vertexCount = static_cast(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(std::min(BATCH, count - first)); + Check(oroMemcpyHtoD(reinterpret_cast(g_Gpu.rays), const_cast(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(function), (batch + BLOCK - 1) / BLOCK, 1, 1, BLOCK, 1, 1, 0, nullptr, arguments, nullptr), + "oroModuleLaunchKernel"); + Check(oroDeviceSynchronize(), "oroDeviceSynchronize"); + Check(oroMemcpyDtoH(static_cast(results) + first * resultBytes, reinterpret_cast(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 Make(const UgcModel::Mesh& mesh) { + if (!Available()) return nullptr; + try { + return std::make_unique(mesh); + } catch (const std::exception&) { + return nullptr; + } + } + + void SetDevice(int index) { + g_DeviceIndex = index; + } +} diff --git a/dUgcServer/Render/UgcRaysHiprt.h b/dUgcServer/Render/UgcRaysHiprt.h new file mode 100644 index 000000000..4502adcac --- /dev/null +++ b/dUgcServer/Render/UgcRaysHiprt.h @@ -0,0 +1,25 @@ +#pragma once + +#include +#include + +#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 Make(const UgcModel::Mesh& mesh); + + // Which device (0: the first); before it is first used + void SetDevice(int index); +} diff --git a/dUgcServer/Render/UgcRender.cpp b/dUgcServer/Render/UgcRender.cpp index 1d7de74ca..d621552fb 100644 --- a/dUgcServer/Render/UgcRender.cpp +++ b/dUgcServer/Render/UgcRender.cpp @@ -284,36 +284,60 @@ namespace UgcRender { return h; } }; - std::unordered_map 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 first; + first.reserve(mesh.positions.size()); + std::vector copyOf(mesh.positions.size(), SIZE_MAX); + std::vector 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(std::lround(mesh.positions[v].x * 1000.0f)), static_cast(std::lround(mesh.positions[v].y * 1000.0f)), static_cast(std::lround(mesh.positions[v].z * 1000.0f)) }, { static_cast(std::lround(normal.x * 100.0f)), static_cast(std::lround(normal.y * 100.0f)), static_cast(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((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(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(1, (1u << 18) / count) : 256; + std::vector batch; + std::vector 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((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(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(open) / static_cast(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(open) / static_cast(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(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(1, (1u << 18) / count) : 1024; + std::vector 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(index) * 0x9E3779B9u; - hash ^= hash >> 16; - hash *= 0x85EBCA6Bu; - hash ^= hash >> 13; - const float turn = static_cast(hash >> 8) / 16777216.0f; - const float shift = static_cast((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(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 batch; + std::vector 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(index) * 0x9E3779B9u; + hash ^= hash >> 16; + hash *= 0x85EBCA6Bu; + hash ^= hash >> 13; + const float turn = static_cast(hash >> 8) / 16777216.0f; + const float shift = static_cast((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(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(open) / static_cast(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(open) / static_cast(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); } } diff --git a/dUgcServer/UgcServer.cpp b/dUgcServer/UgcServer.cpp index 592b895ed..9b3f6d964 100644 --- 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("hiprt_device", 0), 0)); settings.ao.enabled = Setting("bake_ao", 1) != 0; settings.ao.distance = Setting("ao_distance", 5.0f); settings.ao.samples = std::clamp(Setting("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(std::thread::hardware_concurrency(), 1); @@ -691,6 +704,7 @@ int main(int argc, char** argv) { processorConfig.threads = threads > 0 ? threads : std::max(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(processor.Busy()); }); diff --git a/resources/ugcconfig.ini b/resources/ugcconfig.ini index c171aab5e..c90867bee 100644 --- 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) diff --git a/tests/dUgcTests/UgcTests.cpp b/tests/dUgcTests/UgcTests.cpp index 9a012514b..e308a54ac 100644 --- 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 diff --git a/thirdparty/CMakeLists.txt b/thirdparty/CMakeLists.txt index ae34080eb..9a35d860a 100644 --- 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()