From 53d616ce9c853de70e44f9d7b0743f5761213bf1 Mon Sep 17 00:00:00 2001 From: "copilot-swe-agent[bot]" <198982749+Copilot@users.noreply.github.com> Date: Sat, 13 Dec 2025 21:33:10 +0000 Subject: [PATCH 1/4] Initial plan From bd005fd81773a110810ec909d57957335323ab1e Mon Sep 17 00:00:00 2001 From: "copilot-swe-agent[bot]" <198982749+Copilot@users.noreply.github.com> Date: Sat, 13 Dec 2025 21:50:38 +0000 Subject: [PATCH 2/4] Implement unwrapSelection method in _EditKernelsService Add a new method that unrolls cell positions from periodic/modulo space to flat Euclidean space. This is done by: 1. Marking cells near refPos as starting points for unwrapping 2. Propagating unwrapping through connected cells using BFS-like algorithm 3. Using atomicCAS to prevent race conditions during parallel execution The unwrapped positions are stored in Cell::shared1 field. Co-authored-by: chrxh <73127001+chrxh@users.noreply.github.com> --- source/EngineGpuKernels/EditKernels.cu | 62 +++++++++++++++++++ source/EngineGpuKernels/EditKernels.cuh | 2 + source/EngineGpuKernels/EditKernelsService.cu | 18 ++++++ .../EngineGpuKernels/EditKernelsService.cuh | 2 + 4 files changed, 84 insertions(+) diff --git a/source/EngineGpuKernels/EditKernels.cu b/source/EngineGpuKernels/EditKernels.cu index 14d9f8107e..bbd564c057 100644 --- a/source/EngineGpuKernels/EditKernels.cu +++ b/source/EngineGpuKernels/EditKernels.cu @@ -553,6 +553,68 @@ __global__ void cudaRolloutSelectionStep(SimulationData data, int* result) } } +__global__ void cudaMarkUnwrapStartCells(SimulationData data, float2 refPos, float radius) +{ + auto const cellPartition = calcAllThreadsPartition(data.objects.cells.getNumEntries()); + + for (int index = cellPartition.startIndex; index <= cellPartition.endIndex; ++index) { + auto const& cell = data.objects.cells.at(index); + + // Initialize all selected cells: tempValue.as_uint64 = 0 means not yet unwrapped + // After unwrapping, tempValue.as_uint64 = 1 means unwrapped, and shared1 holds the unwrapped position + if (cell->selected == 1 || cell->selected == 2) { + cell->tempValue.as_uint64 = 0; + cell->shared1 = cell->pos; + + // Mark cells within radius of refPos as starting points (already unwrapped) + if (data.cellMap.getDistance(refPos, cell->pos) <= radius) { + cell->tempValue.as_uint64 = 1; + cell->shared1 = cell->pos + data.cellMap.getCorrectionIncrement(refPos, cell->pos); + } + } + } +} + +__global__ void cudaUnwrapSelectionStep(SimulationData data, int* result) +{ + auto const cellPartition = calcAllThreadsPartition(data.objects.cells.getNumEntries()); + + for (int index = cellPartition.startIndex; index <= cellPartition.endIndex; ++index) { + auto const& cell = data.objects.cells.at(index); + + // Only process selected cells that have already been unwrapped + if ((cell->selected == 1 || cell->selected == 2) && cell->tempValue.as_uint64 == 1) { + auto currentCell = cell; + + // Propagate unwrapping to connected cells (heuristics similar to cudaRolloutSelectionStep) + for (int i = 0; i < 30; ++i) { + bool found = false; + for (int j = 0; j < currentCell->numConnections; ++j) { + auto connectedCell = currentCell->connections[j].cell; + // Check if connected cell is selected and not yet unwrapped + if ((connectedCell->selected == 1 || connectedCell->selected == 2) && connectedCell->tempValue.as_uint64 == 0) { + // Atomically mark as unwrapped to prevent race conditions + auto oldValue = atomicCAS(reinterpret_cast(&connectedCell->tempValue.as_uint64), 0ull, 1ull); + if (oldValue == 0) { + // Calculate the unwrapped position based on the current cell's unwrapped position + auto delta = connectedCell->pos - currentCell->pos; + data.cellMap.correctDirection(delta); + connectedCell->shared1 = currentCell->shared1 + delta; + currentCell = connectedCell; + found = true; + atomicExch(result, 1); + break; + } + } + } + if (!found) { + break; + } + } + } + } +} + __global__ void cudaApplyForce(SimulationData data, ApplyForceData applyData) { { diff --git a/source/EngineGpuKernels/EditKernels.cuh b/source/EngineGpuKernels/EditKernels.cuh index d83c6dfde2..85972df8b1 100644 --- a/source/EngineGpuKernels/EditKernels.cuh +++ b/source/EngineGpuKernels/EditKernels.cuh @@ -48,6 +48,8 @@ __global__ void cudaSetSelection(AreaSelectionData selectionData, SimulationData __global__ void cudaRemoveSelection(SimulationData data, bool onlyClusterSelection); __global__ void cudaSwapSelection(float2 pos, float radius, SimulationData data); __global__ void cudaRolloutSelectionStep(SimulationData data, int* result); +__global__ void cudaMarkUnwrapStartCells(SimulationData data, float2 refPos, float radius); +__global__ void cudaUnwrapSelectionStep(SimulationData data, int* result); __global__ void cudaApplyForce(SimulationData data, ApplyForceData applyData); __global__ void cudaResetSelectionResult(SelectionResult result); __global__ void cudaCalcCellWithMinimalPosY(SimulationData data, unsigned long long int* minCellPosYAndIndex); diff --git a/source/EngineGpuKernels/EditKernelsService.cu b/source/EngineGpuKernels/EditKernelsService.cu index dca46c1781..266ffc4e52 100644 --- a/source/EngineGpuKernels/EditKernelsService.cu +++ b/source/EngineGpuKernels/EditKernelsService.cu @@ -8,6 +8,7 @@ _EditKernelsService::_EditKernelsService() { auto& memoryManager = CudaMemoryManager::getInstance(); memoryManager.acquireMemory(1, _cudaRolloutResult); + memoryManager.acquireMemory(1, _cudaUnwrapResult); memoryManager.acquireMemory(1, _cudaSwitchResult); memoryManager.acquireMemory(1, _cudaUpdateResult); memoryManager.acquireMemory(1, _cudaRemoveResult); @@ -26,6 +27,7 @@ _EditKernelsService::~_EditKernelsService() { auto& memoryManager = CudaMemoryManager::getInstance(); memoryManager.freeMemory(_cudaRolloutResult); + memoryManager.freeMemory(_cudaUnwrapResult); memoryManager.freeMemory(_cudaSwitchResult); memoryManager.freeMemory(_cudaUpdateResult); memoryManager.freeMemory(_cudaRemoveResult); @@ -294,6 +296,22 @@ void _EditKernelsService::rolloutSelection(CudaSettings const& gpuSettings, Simu } while (1 == copyToHost(_cudaRolloutResult)); } +void _EditKernelsService::unwrapSelection(CudaSettings const& gpuSettings, SimulationData const& data, float2 const& refPos) +{ + // Mark cells near refPos as starting points for unwrapping + // Use a small radius to find the initial cells in the connected component + KERNEL_CALL(cudaMarkUnwrapStartCells, data, refPos, 1.5f); + cudaDeviceSynchronize(); + + // Propagate unwrapping through connected cells until no more changes + do { + setValueToDevice(_cudaUnwrapResult, 0); + KERNEL_CALL(cudaUnwrapSelectionStep, data, _cudaUnwrapResult); + cudaDeviceSynchronize(); + + } while (1 == copyToHost(_cudaUnwrapResult)); +} + void _EditKernelsService::applyCataclysm(CudaSettings const& gpuSettings, SimulationData const& data) { KERNEL_CALL(cudaApplyCataclysm, data); diff --git a/source/EngineGpuKernels/EditKernelsService.cuh b/source/EngineGpuKernels/EditKernelsService.cuh index d1346b9a22..f114d55e32 100644 --- a/source/EngineGpuKernels/EditKernelsService.cuh +++ b/source/EngineGpuKernels/EditKernelsService.cuh @@ -35,6 +35,7 @@ public: void applyForce(CudaSettings const& gpuSettings, SimulationData const& data, ApplyForceData const& applyData); void rolloutSelection(CudaSettings const& gpuSettings, SimulationData const& data); + void unwrapSelection(CudaSettings const& gpuSettings, SimulationData const& data, float2 const& refPos); void applyCataclysm(CudaSettings const& gpuSettings, SimulationData const& data); @@ -43,6 +44,7 @@ private: // Gpu memory int* _cudaRolloutResult; + int* _cudaUnwrapResult; int* _cudaSwitchResult; int* _cudaUpdateResult; int* _cudaRemoveResult; From 9e187087dfd4bf8924b9a043d5822614fb9d11f7 Mon Sep 17 00:00:00 2001 From: "copilot-swe-agent[bot]" <198982749+Copilot@users.noreply.github.com> Date: Sat, 13 Dec 2025 22:34:51 +0000 Subject: [PATCH 3/4] Fix unwrapSelection to handle all edge cases Addressed review feedback: 1. Fixed: No cells within radius - now finds closest cell regardless 2. Fixed: Multiple cells in same component - only one starting point 3. Fixed: Multiple disconnected components - each gets its own start New algorithm: - cudaInitUnwrapSelection: Initialize all selected cells - cudaFindClosestUnwrappedCell: Find closest unprocessed cell to refPos - cudaMarkUnwrapStartCell: Mark single cell as starting point - cudaUnwrapSelectionStep: BFS propagation through component - Repeat for each disconnected component Co-authored-by: chrxh <73127001+chrxh@users.noreply.github.com> --- source/EngineGpuKernels/EditKernels.cu | 40 ++++++++++++++++--- source/EngineGpuKernels/EditKernels.cuh | 4 +- source/EngineGpuKernels/EditKernelsService.cu | 37 +++++++++++++---- 3 files changed, 66 insertions(+), 15 deletions(-) diff --git a/source/EngineGpuKernels/EditKernels.cu b/source/EngineGpuKernels/EditKernels.cu index bbd564c057..c81cc11328 100644 --- a/source/EngineGpuKernels/EditKernels.cu +++ b/source/EngineGpuKernels/EditKernels.cu @@ -553,7 +553,7 @@ __global__ void cudaRolloutSelectionStep(SimulationData data, int* result) } } -__global__ void cudaMarkUnwrapStartCells(SimulationData data, float2 refPos, float radius) +__global__ void cudaInitUnwrapSelection(SimulationData data) { auto const cellPartition = calcAllThreadsPartition(data.objects.cells.getNumEntries()); @@ -565,16 +565,44 @@ __global__ void cudaMarkUnwrapStartCells(SimulationData data, float2 refPos, flo if (cell->selected == 1 || cell->selected == 2) { cell->tempValue.as_uint64 = 0; cell->shared1 = cell->pos; + } + } +} - // Mark cells within radius of refPos as starting points (already unwrapped) - if (data.cellMap.getDistance(refPos, cell->pos) <= radius) { - cell->tempValue.as_uint64 = 1; - cell->shared1 = cell->pos + data.cellMap.getCorrectionIncrement(refPos, cell->pos); - } +__global__ void cudaFindClosestUnwrappedCell(SimulationData data, float2 refPos, unsigned long long int* minDistanceAndIndex) +{ + auto const cellPartition = calcAllThreadsPartition(data.objects.cells.getNumEntries()); + + for (int index = cellPartition.startIndex; index <= cellPartition.endIndex; ++index) { + auto const& cell = data.objects.cells.at(index); + + // Only consider selected cells that are not yet unwrapped + if ((cell->selected == 1 || cell->selected == 2) && cell->tempValue.as_uint64 == 0) { + // Calculate distance to refPos (using periodic boundary) + auto distance = data.cellMap.getDistance(refPos, cell->pos); + // Encode distance (as fixed-point integer) and index into a single 64-bit value + // Use upper 32 bits for distance (as integer * 1000 for precision), lower 32 bits for index + auto distanceInt = static_cast(distance * 1000.0f); + auto encodedValue = (distanceInt << 32) | static_cast(index); + atomicMin(minDistanceAndIndex, encodedValue); } } } +__global__ void cudaMarkUnwrapStartCell(SimulationData data, int cellIndex, float2 refPos) +{ + if (cellIndex < 0 || cellIndex >= data.objects.cells.getNumEntries()) { + return; + } + + auto const& cell = data.objects.cells.at(cellIndex); + if ((cell->selected == 1 || cell->selected == 2) && cell->tempValue.as_uint64 == 0) { + cell->tempValue.as_uint64 = 1; + // Calculate the unwrapped position: shift cell position relative to refPos + cell->shared1 = cell->pos + data.cellMap.getCorrectionIncrement(refPos, cell->pos); + } +} + __global__ void cudaUnwrapSelectionStep(SimulationData data, int* result) { auto const cellPartition = calcAllThreadsPartition(data.objects.cells.getNumEntries()); diff --git a/source/EngineGpuKernels/EditKernels.cuh b/source/EngineGpuKernels/EditKernels.cuh index 85972df8b1..563402ea12 100644 --- a/source/EngineGpuKernels/EditKernels.cuh +++ b/source/EngineGpuKernels/EditKernels.cuh @@ -48,7 +48,9 @@ __global__ void cudaSetSelection(AreaSelectionData selectionData, SimulationData __global__ void cudaRemoveSelection(SimulationData data, bool onlyClusterSelection); __global__ void cudaSwapSelection(float2 pos, float radius, SimulationData data); __global__ void cudaRolloutSelectionStep(SimulationData data, int* result); -__global__ void cudaMarkUnwrapStartCells(SimulationData data, float2 refPos, float radius); +__global__ void cudaInitUnwrapSelection(SimulationData data); +__global__ void cudaFindClosestUnwrappedCell(SimulationData data, float2 refPos, unsigned long long int* minDistanceAndIndex); +__global__ void cudaMarkUnwrapStartCell(SimulationData data, int cellIndex, float2 refPos); __global__ void cudaUnwrapSelectionStep(SimulationData data, int* result); __global__ void cudaApplyForce(SimulationData data, ApplyForceData applyData); __global__ void cudaResetSelectionResult(SelectionResult result); diff --git a/source/EngineGpuKernels/EditKernelsService.cu b/source/EngineGpuKernels/EditKernelsService.cu index 266ffc4e52..ec3ae97bce 100644 --- a/source/EngineGpuKernels/EditKernelsService.cu +++ b/source/EngineGpuKernels/EditKernelsService.cu @@ -298,18 +298,39 @@ void _EditKernelsService::rolloutSelection(CudaSettings const& gpuSettings, Simu void _EditKernelsService::unwrapSelection(CudaSettings const& gpuSettings, SimulationData const& data, float2 const& refPos) { - // Mark cells near refPos as starting points for unwrapping - // Use a small radius to find the initial cells in the connected component - KERNEL_CALL(cudaMarkUnwrapStartCells, data, refPos, 1.5f); + // Step 1: Initialize all selected cells for unwrapping + KERNEL_CALL(cudaInitUnwrapSelection, data); cudaDeviceSynchronize(); - // Propagate unwrapping through connected cells until no more changes - do { - setValueToDevice(_cudaUnwrapResult, 0); - KERNEL_CALL(cudaUnwrapSelectionStep, data, _cudaUnwrapResult); + // Step 2: Process all connected components + // For each iteration: find the closest unprocessed cell, mark it as start, propagate through its component + bool hasMoreComponents = true; + while (hasMoreComponents) { + // Find the closest unprocessed selected cell to refPos + setValueToDevice(_cudaMinCellPosYAndIndex, 0xffffffffffffffffull); + KERNEL_CALL(cudaFindClosestUnwrappedCell, data, refPos, _cudaMinCellPosYAndIndex); cudaDeviceSynchronize(); - } while (1 == copyToHost(_cudaUnwrapResult)); + auto minDistanceAndIndex = copyToHost(_cudaMinCellPosYAndIndex); + if (minDistanceAndIndex == 0xffffffffffffffffull) { + // No more unprocessed cells + hasMoreComponents = false; + } else { + // Extract cell index from the encoded value (lower 32 bits) + auto closestCellIndex = static_cast(minDistanceAndIndex & 0xffffffff); + + // Mark this cell as the starting point for unwrapping + KERNEL_CALL_1_1(cudaMarkUnwrapStartCell, data, closestCellIndex, refPos); + cudaDeviceSynchronize(); + + // Propagate unwrapping through connected cells in this component + do { + setValueToDevice(_cudaUnwrapResult, 0); + KERNEL_CALL(cudaUnwrapSelectionStep, data, _cudaUnwrapResult); + cudaDeviceSynchronize(); + } while (1 == copyToHost(_cudaUnwrapResult)); + } + } } void _EditKernelsService::applyCataclysm(CudaSettings const& gpuSettings, SimulationData const& data) From 8cebaf4fb60815737d2b21037f01e11a1ee8b781 Mon Sep 17 00:00:00 2001 From: "copilot-swe-agent[bot]" <198982749+Copilot@users.noreply.github.com> Date: Sun, 14 Dec 2025 07:42:10 +0000 Subject: [PATCH 4/4] Optimize unwrapSelection with cluster-based parallel approach Refactored algorithm to process all connected components in parallel: 1. Initialize clusterIndex for each selected cell 2. Find clusters using atomicMin propagation (ClusterProcessor pattern) 3. Find nearest cell to refPos per cluster using atomicMin 4. Mark all starting cells in parallel 5. Propagate unwrapping through all components simultaneously This improves performance from O(components * iterations) to O(iterations) kernel calls when many disconnected components exist. Co-authored-by: chrxh <73127001+chrxh@users.noreply.github.com> --- source/EngineGpuKernels/EditKernels.cu | 102 ++++++++++++++---- source/EngineGpuKernels/EditKernels.cuh | 5 +- source/EngineGpuKernels/EditKernelsService.cu | 44 ++++---- 3 files changed, 101 insertions(+), 50 deletions(-) diff --git a/source/EngineGpuKernels/EditKernels.cu b/source/EngineGpuKernels/EditKernels.cu index c81cc11328..74c256d660 100644 --- a/source/EngineGpuKernels/EditKernels.cu +++ b/source/EngineGpuKernels/EditKernels.cu @@ -560,46 +560,104 @@ __global__ void cudaInitUnwrapSelection(SimulationData data) for (int index = cellPartition.startIndex; index <= cellPartition.endIndex; ++index) { auto const& cell = data.objects.cells.at(index); - // Initialize all selected cells: tempValue.as_uint64 = 0 means not yet unwrapped - // After unwrapping, tempValue.as_uint64 = 1 means unwrapped, and shared1 holds the unwrapped position + // Initialize all selected cells for cluster-based unwrapping + // clusterIndex: used to identify connected components (initialized to self) + // tempValue.as_uint64: 0 = not yet unwrapped, 1 = unwrapped + // shared1: will hold the unwrapped position + // shared2: stores (distance, nearestCellIndex) for cluster heads if (cell->selected == 1 || cell->selected == 2) { - cell->tempValue.as_uint64 = 0; + cell->clusterIndex = index; + cell->tempValue.as_uint64 = 0; // Not unwrapped yet cell->shared1 = cell->pos; + // Use reinterpret_cast to store uint64 in shared2 (two floats = 8 bytes = 64 bits) + *reinterpret_cast(&cell->shared2) = 0xffffffffffffffffull; // Max distance initially } } } -__global__ void cudaFindClosestUnwrappedCell(SimulationData data, float2 refPos, unsigned long long int* minDistanceAndIndex) +__global__ void cudaFindUnwrapClusters(SimulationData data, int* result) { auto const cellPartition = calcAllThreadsPartition(data.objects.cells.getNumEntries()); for (int index = cellPartition.startIndex; index <= cellPartition.endIndex; ++index) { - auto const& cell = data.objects.cells.at(index); + auto currentCell = data.objects.cells.at(index); + + // Only process selected cells + if (currentCell->selected != 1 && currentCell->selected != 2) { + continue; + } - // Only consider selected cells that are not yet unwrapped - if ((cell->selected == 1 || cell->selected == 2) && cell->tempValue.as_uint64 == 0) { - // Calculate distance to refPos (using periodic boundary) - auto distance = data.cellMap.getDistance(refPos, cell->pos); - // Encode distance (as fixed-point integer) and index into a single 64-bit value - // Use upper 32 bits for distance (as integer * 1000 for precision), lower 32 bits for index - auto distanceInt = static_cast(distance * 1000.0f); - auto encodedValue = (distanceInt << 32) | static_cast(index); - atomicMin(minDistanceAndIndex, encodedValue); + // Propagate minimum clusterIndex through connected selected cells (similar to ClusterProcessor) + for (int i = 0; i < 30; ++i) { + bool found = false; + for (int j = 0; j < currentCell->numConnections; ++j) { + auto candidateCell = currentCell->connections[j].cell; + // Only consider connected selected cells + if (candidateCell->selected != 1 && candidateCell->selected != 2) { + continue; + } + auto cellTag = currentCell->clusterIndex; + auto origTag = atomicMin(&candidateCell->clusterIndex, cellTag); + if (cellTag < origTag) { + currentCell = candidateCell; + found = true; + atomicExch(result, 1); + break; + } + } + if (!found) { + break; + } } } } -__global__ void cudaMarkUnwrapStartCell(SimulationData data, int cellIndex, float2 refPos) +__global__ void cudaFindNearestInCluster(SimulationData data, float2 refPos) { - if (cellIndex < 0 || cellIndex >= data.objects.cells.getNumEntries()) { - return; + auto const cellPartition = calcAllThreadsPartition(data.objects.cells.getNumEntries()); + + for (int index = cellPartition.startIndex; index <= cellPartition.endIndex; ++index) { + auto const& cell = data.objects.cells.at(index); + + // Only process selected cells + if (cell->selected != 1 && cell->selected != 2) { + continue; + } + + // Calculate distance to refPos (using periodic boundary) + auto distance = data.cellMap.getDistance(refPos, cell->pos); + + // Encode distance and index into 64-bit value (upper 32 bits: distance * 1000, lower 32 bits: index) + auto distanceInt = static_cast(distance * 1000.0f); + auto encodedValue = (distanceInt << 32) | static_cast(index); + + // Get the cluster head cell and atomically update if this cell is closer + // We store the encoded value in shared2 (reinterpreted as uint64) + auto clusterHead = data.objects.cells.at(cell->clusterIndex); + atomicMin(reinterpret_cast(&clusterHead->shared2), encodedValue); } +} + +__global__ void cudaMarkUnwrapStartCells(SimulationData data, float2 refPos) +{ + auto const cellPartition = calcAllThreadsPartition(data.objects.cells.getNumEntries()); + + for (int index = cellPartition.startIndex; index <= cellPartition.endIndex; ++index) { + auto const& cell = data.objects.cells.at(index); + + // Only process selected cells that are cluster heads (clusterIndex == self) + if ((cell->selected != 1 && cell->selected != 2) || cell->clusterIndex != static_cast(index)) { + continue; + } + + // This is a cluster head - get the nearest cell index from shared2 + auto encodedValue = *reinterpret_cast(&cell->shared2); + auto nearestCellIndex = static_cast(encodedValue & 0xffffffff); - auto const& cell = data.objects.cells.at(cellIndex); - if ((cell->selected == 1 || cell->selected == 2) && cell->tempValue.as_uint64 == 0) { - cell->tempValue.as_uint64 = 1; - // Calculate the unwrapped position: shift cell position relative to refPos - cell->shared1 = cell->pos + data.cellMap.getCorrectionIncrement(refPos, cell->pos); + // Mark the nearest cell as unwrapped with position relative to refPos + auto nearestCell = data.objects.cells.at(nearestCellIndex); + nearestCell->tempValue.as_uint64 = 1; // Mark as unwrapped + nearestCell->shared1 = nearestCell->pos + data.cellMap.getCorrectionIncrement(refPos, nearestCell->pos); } } diff --git a/source/EngineGpuKernels/EditKernels.cuh b/source/EngineGpuKernels/EditKernels.cuh index 563402ea12..d48c5b07b4 100644 --- a/source/EngineGpuKernels/EditKernels.cuh +++ b/source/EngineGpuKernels/EditKernels.cuh @@ -49,8 +49,9 @@ __global__ void cudaRemoveSelection(SimulationData data, bool onlyClusterSelecti __global__ void cudaSwapSelection(float2 pos, float radius, SimulationData data); __global__ void cudaRolloutSelectionStep(SimulationData data, int* result); __global__ void cudaInitUnwrapSelection(SimulationData data); -__global__ void cudaFindClosestUnwrappedCell(SimulationData data, float2 refPos, unsigned long long int* minDistanceAndIndex); -__global__ void cudaMarkUnwrapStartCell(SimulationData data, int cellIndex, float2 refPos); +__global__ void cudaFindUnwrapClusters(SimulationData data, int* result); +__global__ void cudaFindNearestInCluster(SimulationData data, float2 refPos); +__global__ void cudaMarkUnwrapStartCells(SimulationData data, float2 refPos); __global__ void cudaUnwrapSelectionStep(SimulationData data, int* result); __global__ void cudaApplyForce(SimulationData data, ApplyForceData applyData); __global__ void cudaResetSelectionResult(SelectionResult result); diff --git a/source/EngineGpuKernels/EditKernelsService.cu b/source/EngineGpuKernels/EditKernelsService.cu index ec3ae97bce..f95d497db3 100644 --- a/source/EngineGpuKernels/EditKernelsService.cu +++ b/source/EngineGpuKernels/EditKernelsService.cu @@ -298,39 +298,31 @@ void _EditKernelsService::rolloutSelection(CudaSettings const& gpuSettings, Simu void _EditKernelsService::unwrapSelection(CudaSettings const& gpuSettings, SimulationData const& data, float2 const& refPos) { - // Step 1: Initialize all selected cells for unwrapping + // Step 1: Initialize all selected cells (clusterIndex, tempValue, shared1, shared2) KERNEL_CALL(cudaInitUnwrapSelection, data); cudaDeviceSynchronize(); - // Step 2: Process all connected components - // For each iteration: find the closest unprocessed cell, mark it as start, propagate through its component - bool hasMoreComponents = true; - while (hasMoreComponents) { - // Find the closest unprocessed selected cell to refPos - setValueToDevice(_cudaMinCellPosYAndIndex, 0xffffffffffffffffull); - KERNEL_CALL(cudaFindClosestUnwrappedCell, data, refPos, _cudaMinCellPosYAndIndex); + // Step 2: Find connected components among selected cells using cluster propagation + do { + setValueToDevice(_cudaUnwrapResult, 0); + KERNEL_CALL(cudaFindUnwrapClusters, data, _cudaUnwrapResult); cudaDeviceSynchronize(); + } while (1 == copyToHost(_cudaUnwrapResult)); - auto minDistanceAndIndex = copyToHost(_cudaMinCellPosYAndIndex); - if (minDistanceAndIndex == 0xffffffffffffffffull) { - // No more unprocessed cells - hasMoreComponents = false; - } else { - // Extract cell index from the encoded value (lower 32 bits) - auto closestCellIndex = static_cast(minDistanceAndIndex & 0xffffffff); + // Step 3: Find the nearest cell to refPos in each cluster + KERNEL_CALL(cudaFindNearestInCluster, data, refPos); + cudaDeviceSynchronize(); - // Mark this cell as the starting point for unwrapping - KERNEL_CALL_1_1(cudaMarkUnwrapStartCell, data, closestCellIndex, refPos); - cudaDeviceSynchronize(); + // Step 4: Mark the starting cells (nearest in each cluster) as unwrapped + KERNEL_CALL(cudaMarkUnwrapStartCells, data, refPos); + cudaDeviceSynchronize(); - // Propagate unwrapping through connected cells in this component - do { - setValueToDevice(_cudaUnwrapResult, 0); - KERNEL_CALL(cudaUnwrapSelectionStep, data, _cudaUnwrapResult); - cudaDeviceSynchronize(); - } while (1 == copyToHost(_cudaUnwrapResult)); - } - } + // Step 5: Propagate unwrapping through connected cells in all components + do { + setValueToDevice(_cudaUnwrapResult, 0); + KERNEL_CALL(cudaUnwrapSelectionStep, data, _cudaUnwrapResult); + cudaDeviceSynchronize(); + } while (1 == copyToHost(_cudaUnwrapResult)); } void _EditKernelsService::applyCataclysm(CudaSettings const& gpuSettings, SimulationData const& data)