diff --git a/cmake/module_manifest.cmake b/cmake/module_manifest.cmake index 46ff4630..a92b1e17 100644 --- a/cmake/module_manifest.cmake +++ b/cmake/module_manifest.cmake @@ -300,7 +300,7 @@ eve_declare_module(NAME fluids LAYER 5 SCRIPT Fluids SLOT fluids DEPS gpgpu graphics GROUP 3d web) eve_declare_module(NAME procgen LAYER 5 SCRIPT Procgen SLOT procgen - DEPS graphics image map + DEPS gpgpu graphics image map GROUP 3d) eve_declare_module(NAME avatar LAYER 5 SCRIPT Avatar SLOT avatar DEPS animation graphics model3d scene) diff --git a/docs/usr/modules/procgen.md b/docs/usr/modules/procgen.md index 3c303f95..9c0cd3fc 100644 --- a/docs/usr/modules/procgen.md +++ b/docs/usr/modules/procgen.md @@ -361,6 +361,18 @@ metric 或取消状态；并行 Cell/后台任务应各自实例化并重新绑 `getMetricAverageDensity()` 显示密度热度；这些摘要在首次执行时计算并随节点 cache 保存， cache hit 不会再次遍历大型点集。 +### Vulkan / WebGPU Point Compute + +`ProcgenPointGraph.setComputePolicy("auto" | "gpu" | "cpu")` 控制点处理执行器。默认 `auto` +在输入达到 `setComputeMinimumPoints()`（默认 1024）后，将 `transform` 的位置、法线、旋转和 +缩放批量打包到 storage buffer，通过单个 Sequence 完成上传、dispatch 和回读；Vulkan 使用 +GLSL/SPIR-V，WebGPU 使用 WGSL。设备尚未初始化、shader 编译、提交或回读失败时会确定性地 +回退现有 CPU 实现，`getComputeFallbackReason()` 返回原因。 + +`getMetricBackend(index)` 返回该节点实际使用的 `cpu`、`vulkan` 或 `webgpu`，`debugReport()` +也包含 `backend=`。这使编辑器和性能采集能够区分 GPU 命中与 CPU fallback。当前 GPU kernel +覆盖 `transform`；其余节点继续走 CPU，后续可在同一执行层扩展为连续节点融合，无需改变图资产。 + 子图节点将首个输入写入嵌套图指定的 input node，并返回指定 output node： ```nut @@ -792,7 +804,7 @@ local castle = gen.generateMesh("mesh.castle", p, gfx); - `getDensity()`、`getError()`、`getFloatAttribute()`、`getNormalX()`、`getNormalY()`、`getNormalZ()`、`getOutput()`、`getOutputCount()`、`getOutputName()`、`getPointSeed()`、`getScaleX()`、`getScaleY()`、`getScaleZ()`、`getStringAttribute()`、`getSystemDebugReport()`、`getSystemOutput()`、`getSystemOutputCount()`、`getSystemOutputName()`、`getSystemRevision()`、`getSystemSeed()`、`getTraceCount()`、`getTraceInputCount()`、`getTraceMilliseconds()`、`getTraceName()`、`getTraceOutputCount()`、`getTriangleGroup()`、`getX()`、`getY()`、`getYaw()`、`getZ()`、`hasFailed()`、`hasFloatAttribute()`、`hasOutput()`、`hasStringAttribute()`、`hasSystem()`、`isActive()`、`jitterPoints()`、`lastError()`、`newGrid()`、`newOutput()`、`newParams()`、`newPointSet()`、`publish()`、`randomSeed()`、`removeSystem()`、`resize()`、`sampleGrid()`、`seedFor()`、`selfPrune()`、`setCell()`、`setDensity()`、`setDetail()`、`setFloat()`、`setFloatAttribute()`、`setInt()` - `setBool()`、`setLayer()`、`setMeta()`、`setNormal()`、`setPalette()`、`setPaletteGid()`、`setPath()`、`setPointSeed()`、`setPosition()`、`setScale()`、`setSeed()`、`setSize()`、`setString()`、`setStringAttribute()`、`setYaw()`、`trace()`、`uploadMesh()`、`setActiveGroup()` - `setTarget()` -- UE PCG 扩展：`clearCache()`、`clearGenerationSources()`、`clearParameterOverride()`、`disconnect()`、`exposeParameter()`、`getActiveCellCount()`、`getAssetCount()`、`getAssetName()`、`getCacheHitCount()`、`getCellRevision()`、`getDirectionWeight()`、`getExclusionCount()`、`getExecutionCount()`、`getExecutionNodeBudget()`、`getFailedCellCount()`、`getFrameTimeBudget()`、`getFrustumBehindRadius()`、`getFrustumHalfAngle()`、`getGeneratingCount()`、`getGenerationSourceCount()`、`getGenerationSourceId()`、`getInputNode()`、`getLayerCount()`、`getLayerDensity()`、`getLayerName()`、`getLayerPriority()`、`getLevelCellSize()`、`getLevelCleanupRadius()`、`getLevelCount()`、`getLevelGenerationRadius()`、`getMaxActiveCells()`、`getMaxGenerating()`、`getMaxGenerationRetries()`、`getMaxY()`、`getMetricCount()`、`getMetricMilliseconds()`、`getMetricNodeId()`、`getMetricOutputCount()`、`getMinY()`、`getModuleCount()`、`getModuleSymbol()`、`getNodeCount()`、`getNodeId()`、`getNodeOperation()`、`getOperationInputCount()`、`getOperationParamCount()`、`getOperationParamDefault()`、`getOperationParamKey()`、`getOperationParamKind()`、`getParameterCount()`、`getParameterFloat()`、`getParameterInt()`、`getParameterKind()`、`getParameterName()`、`getParameterString()`、`getPendingCleanupCount()`、`getPendingGenerateCount()`、`getRevision()`、`getTicket()`、`getVariantAsset()`、`getVariantCount()`、`getVariantLength()`、`hasCell()`、`hasLayer()`、`hasModule()`、`hasNode()`、`hasParameterOverride()`、`intersectSpatial()`、`isFrustumCullingEnabled()`、`isMetricCacheHit()`、`pointData()`、`refreshGenerationSources()`、`removeLayer()`、`removeModule()`、`removeNode()`、`requestCancel()`、`resetCancellation()`、`retryFailedCells()`、`setExecutionNodeBudget()`、`setMaxActiveCells()`、`setMaxGenerationRetries()`、`setNodeString()`、`setParameterFloat()`、`setParameterInt()`、`setParameterString()`、`wasCancelled()`。 +- UE PCG 扩展：`clearCache()`、`clearGenerationSources()`、`clearParameterOverride()`、`disconnect()`、`exposeParameter()`、`getActiveCellCount()`、`getAssetCount()`、`getAssetName()`、`getCacheHitCount()`、`getCellRevision()`、`getComputeFallbackReason()`、`getComputeMinimumPoints()`、`getComputePolicy()`、`getDirectionWeight()`、`getExclusionCount()`、`getExecutionCount()`、`getExecutionNodeBudget()`、`getFailedCellCount()`、`getFrameTimeBudget()`、`getFrustumBehindRadius()`、`getFrustumHalfAngle()`、`getGeneratingCount()`、`getGenerationSourceCount()`、`getGenerationSourceId()`、`getInputNode()`、`getLayerCount()`、`getLayerDensity()`、`getLayerName()`、`getLayerPriority()`、`getLevelCellSize()`、`getLevelCleanupRadius()`、`getLevelCount()`、`getLevelGenerationRadius()`、`getMaxActiveCells()`、`getMaxGenerating()`、`getMaxGenerationRetries()`、`getMaxY()`、`getMetricBackend()`、`getMetricCount()`、`getMetricMilliseconds()`、`getMetricNodeId()`、`getMetricOutputCount()`、`getMinY()`、`getModuleCount()`、`getModuleSymbol()`、`getNodeCount()`、`getNodeId()`、`getNodeOperation()`、`getOperationInputCount()`、`getOperationParamCount()`、`getOperationParamDefault()`、`getOperationParamKey()`、`getOperationParamKind()`、`getParameterCount()`、`getParameterFloat()`、`getParameterInt()`、`getParameterKind()`、`getParameterName()`、`getParameterString()`、`getPendingCleanupCount()`、`getPendingGenerateCount()`、`getRevision()`、`getTicket()`、`getVariantAsset()`、`getVariantCount()`、`getVariantLength()`、`hasCell()`、`hasLayer()`、`hasModule()`、`hasNode()`、`hasParameterOverride()`、`intersectSpatial()`、`isFrustumCullingEnabled()`、`isMetricCacheHit()`、`pointData()`、`refreshGenerationSources()`、`removeLayer()`、`removeModule()`、`removeNode()`、`requestCancel()`、`resetCancellation()`、`retryFailedCells()`、`setComputeMinimumPoints()`、`setComputePolicy()`、`setExecutionNodeBudget()`、`setMaxActiveCells()`、`setMaxGenerationRetries()`、`setNodeString()`、`setParameterFloat()`、`setParameterInt()`、`setParameterString()`、`wasCancelled()`。 ## 使用要点 diff --git a/src/modules/procgen/PointCompute.cpp b/src/modules/procgen/PointCompute.cpp new file mode 100644 index 00000000..d71f2b76 --- /dev/null +++ b/src/modules/procgen/PointCompute.cpp @@ -0,0 +1,210 @@ +#include "procgen/PointCompute.h" + +#include "common/Module.h" +#include "gpgpu/ComputeShader.h" +#include "gpgpu/Gpgpu.h" +#include "gpgpu/GpuBuffer.h" +#include "gpgpu/Sequence.h" + +#include +#include +#include +#include +#include +#include + +namespace eve::procgen { +namespace { + +constexpr int kPointFloats = 11; +constexpr int kWorkgroupSize = 64; + +const char* transformKernel() { +#ifdef EVENGINE_WEBGPU + return R"( +struct Data { values: array }; +struct Push { data: array }; +@group(0) @binding(0) var points: Data; +@group(0) @binding(8) var push: Push; +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3u) { + if (gid.x >= u32(push.data[0].x)) { return; } + let base = gid.x * 11u; + let sx = push.data[1].x; + let sy = push.data[1].y; + let sz = push.data[1].z; + let c = push.data[2].x; + let s = push.data[2].y; + let x = points.values[base] * sx; + let z = points.values[base + 2u] * sz; + points.values[base] = x * c - z * s + push.data[0].y; + points.values[base + 1u] = points.values[base + 1u] * sy + push.data[0].z; + points.values[base + 2u] = x * s + z * c + push.data[0].w; + var nx = points.values[base + 3u] / select(1.0, sx, abs(sx) > 0.000001); + var ny = points.values[base + 4u] / select(1.0, sy, abs(sy) > 0.000001); + var nz = points.values[base + 5u] / select(1.0, sz, abs(sz) > 0.000001); + let rotated = vec3f(nx * c - nz * s, ny, nx * s + nz * c); + let lengthSquared = dot(rotated, rotated); + let normal = select(rotated, normalize(rotated), lengthSquared > 0.0); + points.values[base + 3u] = normal.x; + points.values[base + 4u] = normal.y; + points.values[base + 5u] = normal.z; + points.values[base + 6u] += push.data[2].z; + points.values[base + 7u] *= sx; + points.values[base + 8u] *= sy; + points.values[base + 9u] *= sz; +} +)"; +#else + return R"(#version 450 +layout(local_size_x = 64) in; +layout(set = 0, binding = 0) buffer Data { float values[]; } points; +layout(push_constant) uniform Push { float data[32]; } push; +void main() { + uint i = gl_GlobalInvocationID.x; + if (i >= uint(push.data[0])) return; + uint base = i * 11u; + float sx = push.data[4], sy = push.data[5], sz = push.data[6]; + float c = push.data[8], s = push.data[9]; + float x = points.values[base] * sx; + float z = points.values[base + 2u] * sz; + points.values[base] = x * c - z * s + push.data[1]; + points.values[base + 1u] = points.values[base + 1u] * sy + push.data[2]; + points.values[base + 2u] = x * s + z * c + push.data[3]; + float nx = points.values[base + 3u] / (abs(sx) > 0.000001 ? sx : 1.0); + float ny = points.values[base + 4u] / (abs(sy) > 0.000001 ? sy : 1.0); + float nz = points.values[base + 5u] / (abs(sz) > 0.000001 ? sz : 1.0); + vec3 normal = vec3(nx * c - nz * s, ny, nx * s + nz * c); + float normalLength = length(normal); + if (normalLength > 0.0) normal /= normalLength; + points.values[base + 3u] = normal.x; + points.values[base + 4u] = normal.y; + points.values[base + 5u] = normal.z; + points.values[base + 6u] += push.data[10]; + points.values[base + 7u] *= sx; + points.values[base + 8u] *= sy; + points.values[base + 9u] *= sz; +} +)"; +#endif +} + +} // namespace + +struct PointCompute::Impl { + std::unique_ptr storage; + std::unique_ptr staging; + std::unique_ptr shader; + std::unique_ptr sequence; + int capacityBytes = 0; + + void reset() { + sequence.reset(); + shader.reset(); + staging.reset(); + storage.reset(); + capacityBytes = 0; + } +}; + +PointCompute::PointCompute() : impl_(std::make_unique()) {} +PointCompute::~PointCompute() = default; + +bool PointCompute::transform(const PointSet& input, PointSet& output, float translateX, float translateY, + float translateZ, float yawDegrees, float scaleX, float scaleY, float scaleZ) { + error_.clear(); + if (input.empty()) { + output = input; + return true; + } + constexpr size_t bytesPerPoint = size_t(kPointFloats) * sizeof(float); + if (size_t(input.getCount()) > size_t(std::numeric_limits::max()) / bytesPerPoint) { + error_ = "point data exceeds GPU buffer address range"; + return false; + } + try { + auto* gpu = eve::ModuleManager::getInstance("Gpgpu"); + if (!gpu) gpu = eve::gpgpu::Gpgpu::create(); + if (!gpu || !gpu->isAvailable()) { + error_ = "compute device unavailable"; + return false; + } + + std::vector packed(size_t(input.getCount()) * kPointFloats); + for (int index = 0; index < input.getCount(); ++index) { + const ProcgenPoint& point = input.points()[size_t(index)]; + float* values = packed.data() + size_t(index) * kPointFloats; + values[0] = point.x; + values[1] = point.y; + values[2] = point.z; + values[3] = point.normalX; + values[4] = point.normalY; + values[5] = point.normalZ; + values[6] = point.yaw; + values[7] = point.scaleX; + values[8] = point.scaleY; + values[9] = point.scaleZ; + values[10] = point.density; + } + const int byteSize = int(packed.size() * sizeof(float)); + if (!impl_->shader) impl_->shader.reset(gpu->newShader(transformKernel())); + if (!impl_->sequence) impl_->sequence.reset(gpu->newSequence()); + if (byteSize > impl_->capacityBytes) { + int capacity = std::max(byteSize, impl_->capacityBytes + impl_->capacityBytes / 2); + impl_->storage.reset(gpu->newBuffer(capacity, "storage")); + impl_->staging.reset(gpu->newBuffer(capacity, "staging")); + impl_->capacityBytes = capacity; + } + if (!impl_->storage || !impl_->staging || !impl_->shader || !impl_->sequence || + !impl_->sequence->isAvailable()) { + error_ = "compute resources unavailable"; + return false; + } + impl_->shader->bindBuffer(0, impl_->storage.get()); + constexpr float degreesToRadians = 0.017453292519943295f; + const float radians = yawDegrees * degreesToRadians; + impl_->shader->setFloat(0, float(input.getCount())); + impl_->shader->setFloat(1, translateX); + impl_->shader->setFloat(2, translateY); + impl_->shader->setFloat(3, translateZ); + impl_->shader->setFloat(4, scaleX); + impl_->shader->setFloat(5, scaleY); + impl_->shader->setFloat(6, scaleZ); + impl_->shader->setFloat(8, std::cos(radians)); + impl_->shader->setFloat(9, std::sin(radians)); + impl_->shader->setFloat(10, yawDegrees); + impl_->sequence->begin(); + impl_->sequence->recordUpload(impl_->storage.get(), packed.data(), uint64_t(byteSize)); + impl_->sequence->recordDispatch(impl_->shader.get(), (input.getCount() + kWorkgroupSize - 1) / kWorkgroupSize); + impl_->sequence->recordDownload(impl_->storage.get(), impl_->staging.get(), uint64_t(byteSize)); + impl_->sequence->submit(); + impl_->staging->downloadBytes(packed.data(), uint64_t(byteSize)); + + output = input; + for (int index = 0; index < output.getCount(); ++index) { + ProcgenPoint& point = output.points()[size_t(index)]; + const float* values = packed.data() + size_t(index) * kPointFloats; + point.x = values[0]; + point.y = values[1]; + point.z = values[2]; + point.normalX = values[3]; + point.normalY = values[4]; + point.normalZ = values[5]; + point.yaw = values[6]; + point.scaleX = values[7]; + point.scaleY = values[8]; + point.scaleZ = values[9]; + point.density = values[10]; + } + return true; + } catch (const std::exception& exception) { + error_ = exception.what(); + impl_->reset(); + } catch (...) { + error_ = "unknown compute failure"; + impl_->reset(); + } + return false; +} + +} // namespace eve::procgen diff --git a/src/modules/procgen/PointCompute.h b/src/modules/procgen/PointCompute.h new file mode 100644 index 00000000..b44bc911 --- /dev/null +++ b/src/modules/procgen/PointCompute.h @@ -0,0 +1,35 @@ +#pragma once + +#include "procgen/PointSet.h" + +#include +#include + +namespace eve::procgen { + +/** @brief Backend-neutral GPU kernels used by PointGraph with deterministic CPU fallback. */ +class PointCompute { +public: + PointCompute(); + ~PointCompute(); + PointCompute(const PointCompute&) = delete; + PointCompute& operator=(const PointCompute&) = delete; + + /** + * @brief Transform attributed points through the active Vulkan or WebGPU compute backend. + * @return True when GPU execution completed; false leaves output untouched and describes + * the fallback reason in getError(). + */ + bool transform(const PointSet& input, PointSet& output, float translateX, float translateY, float translateZ, + float yawDegrees, float scaleX, float scaleY, float scaleZ); + + /** @brief Return the last GPU unavailability, compilation, dispatch, or readback error. */ + const std::string& getError() const { return error_; } + +private: + struct Impl; + std::unique_ptr impl_; + std::string error_; +}; + +} // namespace eve::procgen diff --git a/src/modules/procgen/PointGraph.cpp b/src/modules/procgen/PointGraph.cpp index 137a8c0a..520672ab 100644 --- a/src/modules/procgen/PointGraph.cpp +++ b/src/modules/procgen/PointGraph.cpp @@ -1,6 +1,7 @@ #include "procgen/PointGraph.h" #include "procgen/Biome.h" +#include "procgen/PointCompute.h" #include "procgen/ShapeGrammar.h" #include @@ -88,13 +89,14 @@ uint64_t nowNanoseconds() { .count()); } -PointGraphNodeMetric makeMetric(const std::string& id, const PointSet& points, - float milliseconds, bool cacheHit) { +PointGraphNodeMetric makeMetric(const std::string& id, const PointSet& points, float milliseconds, bool cacheHit, + const std::string& backend = "cpu") { PointGraphNodeMetric metric; metric.id = id; metric.outputCount = points.getCount(); metric.milliseconds = milliseconds; metric.cacheHit = cacheHit; + metric.backend = backend; if (points.empty()) return metric; const auto& first = points.points().front(); metric.minX = metric.maxX = first.x; @@ -133,6 +135,8 @@ uint64_t spatialSampleUpperBound(const SpatialData& spatial, float spacing) { } // namespace +PointGraph::PointGraph() : pointCompute_(std::make_shared()) {} + bool PointGraph::addNode(const std::string& id, const std::string& operation) { if (id.empty() || !operationKnown(operation) || nodes_.find(id) != nodes_.end()) return false; nodes_[id].id = id; @@ -376,6 +380,7 @@ PointSet* PointGraph::execute(const std::string& outputId) { cacheHitCount_ = 0; evaluatedNodes_ = 0; lastCancelled_ = false; + computeFallbackReason_.clear(); ++executionCount_; if (cancelRequested_) { error_ = "execution cancelled"; @@ -396,6 +401,21 @@ void PointGraph::setMaxNodeOutputPoints(int points) { clearCache(); } int PointGraph::getMaxNodeOutputPoints() const { return maxNodeOutputPoints_; } +bool PointGraph::setComputePolicy(const std::string& policy) { + if (policy != "auto" && policy != "gpu" && policy != "cpu") return false; + if (computePolicy_ == policy) return true; + computePolicy_ = policy; + clearCache(); + return true; +} +std::string PointGraph::getComputePolicy() const { return computePolicy_; } +void PointGraph::setComputeMinimumPoints(int points) { + const int clamped = std::max(0, points); + if (computeMinimumPoints_ == clamped) return; + computeMinimumPoints_ = clamped; + clearCache(); +} +int PointGraph::getComputeMinimumPoints() const { return computeMinimumPoints_; } void PointGraph::requestCancel() { cancelRequested_ = true; } void PointGraph::resetCancellation() { cancelRequested_ = false; @@ -432,6 +452,10 @@ float PointGraph::getMetricMilliseconds(int index) const { bool PointGraph::isMetricCacheHit(int index) const { return index >= 0 && index < int(metrics_.size()) && metrics_[size_t(index)].cacheHit; } +std::string PointGraph::getMetricBackend(int index) const { + return index >= 0 && index < int(metrics_.size()) ? metrics_[size_t(index)].backend : std::string(); +} +std::string PointGraph::getComputeFallbackReason() const { return computeFallbackReason_; } float PointGraph::getMetricMinX(int index) const { return index >= 0 && index < int(metrics_.size()) ? metrics_[size_t(index)].minX : 0.f; } @@ -466,8 +490,9 @@ std::string PointGraph::debugReport() const { out << "nodes=" << nodes_.size() << " executions=" << executionCount_ << " cacheHits=" << cacheHitCount_; for (const auto& metric : metrics_) - out << "\n" << metric.id << " count=" << metric.outputCount << " ms=" - << metric.milliseconds << (metric.cacheHit ? " cached" : ""); + out << "\n" + << metric.id << " count=" << metric.outputCount << " ms=" << metric.milliseconds + << " backend=" << metric.backend << (metric.cacheHit ? " cached" : ""); if (!error_.empty()) out << "\nerror=" << error_; return out.str(); } @@ -532,6 +557,7 @@ const PointSet* PointGraph::evaluate(const std::string& id, } states[id] = 1; const uint64_t started = nowNanoseconds(); + std::string nodeBackend = "cpu"; const PointSet* first = nullptr; const PointSet* second = nullptr; if (!node.inputs[0].empty()) { @@ -649,12 +675,30 @@ const PointSet* PointGraph::evaluate(const std::string& id, floatValue(node, "defaultValue", 0.f)); } else if (node.operation == "transform") { if (!first) error_ = "transform requires input: " + id; - else - node.cache = transformPointSet( - *first, floatValue(node, "x", 0.f), floatValue(node, "y", 0.f), - floatValue(node, "z", 0.f), floatValue(node, "yaw", 0.f), - floatValue(node, "scaleX", 1.f), floatValue(node, "scaleY", 1.f), - floatValue(node, "scaleZ", 1.f)); + else { + const bool gpuEligible = + !first->empty() && + (computePolicy_ == "gpu" || (computePolicy_ == "auto" && first->getCount() >= computeMinimumPoints_)); + const bool gpuComplete = + gpuEligible && + pointCompute_->transform(*first, node.cache, floatValue(node, "x", 0.f), floatValue(node, "y", 0.f), + floatValue(node, "z", 0.f), floatValue(node, "yaw", 0.f), + floatValue(node, "scaleX", 1.f), floatValue(node, "scaleY", 1.f), + floatValue(node, "scaleZ", 1.f)); + if (gpuComplete) { +#ifdef EVENGINE_WEBGPU + nodeBackend = "webgpu"; +#else + nodeBackend = "vulkan"; +#endif + } else { + if (gpuEligible) computeFallbackReason_ = pointCompute_->getError(); + node.cache = transformPointSet(*first, floatValue(node, "x", 0.f), floatValue(node, "y", 0.f), + floatValue(node, "z", 0.f), floatValue(node, "yaw", 0.f), + floatValue(node, "scaleX", 1.f), floatValue(node, "scaleY", 1.f), + floatValue(node, "scaleZ", 1.f)); + } + } } else if (node.operation == "filter.float") { if (!first) error_ = "filter.float requires input: " + id; else @@ -710,6 +754,8 @@ const PointSet* PointGraph::evaluate(const std::string& id, if (!first || !node.subgraph) error_ = "subgraph requires input and graph: " + id; else { node.subgraph->setMaxNodeOutputPoints(maxNodeOutputPoints_); + node.subgraph->setComputePolicy(computePolicy_); + node.subgraph->setComputeMinimumPoints(computeMinimumPoints_); if (!node.subgraph->setNodePoints(node.subgraphInput, const_cast(first))) error_ = "subgraph input failed at " + id + ": " + node.subgraph->getError(); else { @@ -731,7 +777,7 @@ const PointSet* PointGraph::evaluate(const std::string& id, node.cacheValid = true; states[id] = 2; const float elapsed = float(nowNanoseconds() - started) * 0.000001f; - node.cacheMetric = makeMetric(id, node.cache, elapsed, false); + node.cacheMetric = makeMetric(id, node.cache, elapsed, false, nodeBackend); metrics_.push_back(node.cacheMetric); return &node.cache; } diff --git a/src/modules/procgen/PointGraph.h b/src/modules/procgen/PointGraph.h index afd0b07b..1f126b45 100644 --- a/src/modules/procgen/PointGraph.h +++ b/src/modules/procgen/PointGraph.h @@ -12,6 +12,7 @@ namespace eve::procgen { class BiomeRules; +class PointCompute; class ShapeGrammar; /** @brief Per-node execution information from the most recent graph run. */ @@ -20,6 +21,7 @@ struct PointGraphNodeMetric { int outputCount = 0; float milliseconds = 0.f; bool cacheHit = false; + std::string backend = "cpu"; float minX = 0.f; float minY = 0.f; float minZ = 0.f; @@ -39,6 +41,7 @@ struct PointGraphNodeMetric { */ class PointGraph { public: + PointGraph(); /** @brief Add a node using one of the operations returned by operationAt(). */ bool addNode(const std::string& id, const std::string& operation); /** @brief Remove a node and every edge that references it. */ @@ -105,7 +108,16 @@ public: int getExecutionNodeBudget() const; /** @brief Limit points produced by any node; zero disables the limit. */ void setMaxNodeOutputPoints(int points); + /** @brief Return the per-node point output limit; zero means unlimited. */ int getMaxNodeOutputPoints() const; + /** @brief Select `auto`, `gpu`, or `cpu` point execution; GPU failures fall back to CPU. */ + bool setComputePolicy(const std::string& policy); + /** @brief Return the configured point compute policy. */ + std::string getComputePolicy() const; + /** @brief Set the minimum point count for `auto` GPU dispatch. */ + void setComputeMinimumPoints(int points); + /** @brief Return the minimum point count for automatic GPU dispatch. */ + int getComputeMinimumPoints() const; /** @brief Cancel subsequent execution until resetCancellation is called. */ void requestCancel(); void resetCancellation(); @@ -124,6 +136,10 @@ public: int getMetricOutputCount(int index) const; float getMetricMilliseconds(int index) const; bool isMetricCacheHit(int index) const; + /** @brief Return `cpu`, `vulkan`, or `webgpu` for one evaluated node. */ + std::string getMetricBackend(int index) const; + /** @brief Return why the latest eligible GPU node used the CPU fallback. */ + std::string getComputeFallbackReason() const; /** @brief Return cached output bounds for editor/debug visualization. */ float getMetricMinX(int index) const; /** @brief Return the cached minimum output Y coordinate. */ @@ -211,6 +227,10 @@ private: uint64_t revision_ = 0; int executionNodeBudget_ = 0; int maxNodeOutputPoints_ = 0; + std::string computePolicy_ = "auto"; + int computeMinimumPoints_ = 1024; + std::string computeFallbackReason_; + std::shared_ptr pointCompute_; int evaluatedNodes_ = 0; bool cancelRequested_ = false; bool lastCancelled_ = false; diff --git a/src/modules/procgen/PointGraphScript.cpp b/src/modules/procgen/PointGraphScript.cpp index 7f6c9cd3..f3cd5506 100644 --- a/src/modules/procgen/PointGraphScript.cpp +++ b/src/modules/procgen/PointGraphScript.cpp @@ -46,6 +46,10 @@ void exposePointGraph(ssq::Table& table) { graph.addFunc("getExecutionNodeBudget", &PointGraph::getExecutionNodeBudget); graph.addFunc("setMaxNodeOutputPoints", &PointGraph::setMaxNodeOutputPoints); graph.addFunc("getMaxNodeOutputPoints", &PointGraph::getMaxNodeOutputPoints); + graph.addFunc("setComputePolicy", &PointGraph::setComputePolicy); + graph.addFunc("getComputePolicy", &PointGraph::getComputePolicy); + graph.addFunc("setComputeMinimumPoints", &PointGraph::setComputeMinimumPoints); + graph.addFunc("getComputeMinimumPoints", &PointGraph::getComputeMinimumPoints); graph.addFunc("requestCancel", &PointGraph::requestCancel); graph.addFunc("resetCancellation", &PointGraph::resetCancellation); graph.addFunc("wasCancelled", &PointGraph::wasCancelled); @@ -60,6 +64,8 @@ void exposePointGraph(ssq::Table& table) { graph.addFunc("getMetricOutputCount", &PointGraph::getMetricOutputCount); graph.addFunc("getMetricMilliseconds", &PointGraph::getMetricMilliseconds); graph.addFunc("isMetricCacheHit", &PointGraph::isMetricCacheHit); + graph.addFunc("getMetricBackend", &PointGraph::getMetricBackend); + graph.addFunc("getComputeFallbackReason", &PointGraph::getComputeFallbackReason); graph.addFunc("getMetricMinX", &PointGraph::getMetricMinX); graph.addFunc("getMetricMinY", &PointGraph::getMetricMinY); graph.addFunc("getMetricMinZ", &PointGraph::getMetricMinZ); diff --git a/test/gpgpu.cpp b/test/gpgpu.cpp index b98209b7..71a5e5a6 100644 --- a/test/gpgpu.cpp +++ b/test/gpgpu.cpp @@ -32,6 +32,7 @@ #include "graphics/Volumetric.h" #include "graphics/Water.h" #include "graphics/Waterfall.h" +#include "procgen/PointCompute.h" #include "window/Window.h" #include @@ -279,6 +280,37 @@ TEST_CASE("gpgpu.sequence.dispatchScale") { delete staging; } +TEST_CASE("gpgpu.procgen.transformParity") { + if (!tryInitHeadlessGfx()) return; + constexpr int count = 2048; + eve::procgen::PointSet input; + for (int index = 0; index < count; ++index) { + const int point = input.add(float(index % 64), float(index % 7), float(index / 64)); + input.setNormal(point, 0.25f, 1.f, -0.5f); + input.setYaw(point, float(index % 360)); + input.setScale(point, 0.5f, 1.5f, 2.f); + input.setDensity(point, float(index % 10) * 0.1f); + input.setPointSeed(point, uint32_t(index + 17)); + input.setStringAttribute(point, "asset", "oak"); + } + const eve::procgen::PointSet expected = + eve::procgen::transformPointSet(input, 12.f, -3.f, 8.f, 31.f, 1.25f, 0.75f, 2.f); + eve::procgen::PointSet actual; + eve::procgen::PointCompute compute; + REQUIRE(compute.transform(input, actual, 12.f, -3.f, 8.f, 31.f, 1.25f, 0.75f, 2.f)); + REQUIRE(actual.getCount() == expected.getCount()); + for (int index = 0; index < count; ++index) { + CHECK(std::fabs(actual.getX(index) - expected.getX(index)) < 0.0001f); + CHECK(std::fabs(actual.getY(index) - expected.getY(index)) < 0.0001f); + CHECK(std::fabs(actual.getZ(index) - expected.getZ(index)) < 0.0001f); + CHECK(std::fabs(actual.getNormalX(index) - expected.getNormalX(index)) < 0.0001f); + CHECK(std::fabs(actual.getNormalY(index) - expected.getNormalY(index)) < 0.0001f); + CHECK(std::fabs(actual.getNormalZ(index) - expected.getNormalZ(index)) < 0.0001f); + CHECK(actual.getPointSeed(index) == input.getPointSeed(index)); + CHECK(actual.getStringAttribute(index, "asset", "") == "oak"); + } +} + TEST_CASE("gpgpu.sequence.singleSubmitChainedDispatches") { if (!tryInitHeadlessGfx()) return; auto *mod = Gpgpu::create(); diff --git a/test/procgen_point_graph_stress.cpp b/test/procgen_point_graph_stress.cpp index 452d051a..70ff7e16 100644 --- a/test/procgen_point_graph_stress.cpp +++ b/test/procgen_point_graph_stress.cpp @@ -58,6 +58,29 @@ TEST_CASE("procgen.pointGraphStress.resumesDeepGraphAcrossBudgetWindows") { CHECK(graph.getCacheHitCount() > 0); } +TEST_CASE("procgen.pointGraphStress.computePolicyPreservesDeterministicCpuPath") { + PointSet input; + input.add(2.f, 3.f, 4.f); + PointGraph graph; + CHECK(graph.addNode("source", "input")); + CHECK(graph.addNode("move", "transform")); + CHECK(graph.connect("source", "move")); + CHECK(graph.setNodePoints("source", &input)); + CHECK(graph.setNodeFloat("move", "x", 5.f)); + CHECK(!graph.setComputePolicy("invalid")); + CHECK(graph.setComputePolicy("cpu")); + graph.setComputeMinimumPoints(0); + + std::unique_ptr output(graph.execute("move")); + REQUIRE(bool(output)); + CHECK_EQ(output->getX(0), 7.f); + CHECK_EQ(graph.getComputePolicy(), "cpu"); + CHECK_EQ(graph.getComputeMinimumPoints(), 0); + CHECK_EQ(graph.getMetricBackend(graph.getMetricCount() - 1), "cpu"); + CHECK(graph.getComputeFallbackReason().empty()); + CHECK(graph.debugReport().find("backend=cpu") != std::string::npos); +} + TEST_CASE("procgen.pointGraphStress.isBitStableAcrossIndependentBiomeInstances") { SpatialData domain = SpatialData::box(-32.f, 0.f, -32.f, 32.f, 0.f, 32.f); SpatialData center = SpatialData::sphere(0.f, 0.f, 0.f, 12.f);