Singleton device allocator#17
Merged
Merged
Conversation
bugfix: Thread aware random number generator for SYCL chod: split out sycl support to separate header.
Build using DPCPP=1 DEBUG=1 The run using "env ONEAPI_DEVICE_SELECTOR=*:cpu python3 spatial_ecolab.py"
bugfix: fatalErrorFlag (SYCL)
…pression should allocate on leader, the copy in parallel
…ted solely ion one thread of work group, and all allocations called with same allocation on all threads, and return the same pointer to all threads. Still crashing after about 3E5 timesteps of spatial_ecolab.py.
|
Warning Review limit reached
Next review available in: 11 minutes Enable usage-based reviews in Billing to review now. Otherwise, wait until the next included review is available. How can I continue?After more reviews become available, a review can be triggered using the To avoid repeated limits, reduce automatic review volume by pausing incremental auto-reviews earlier, using label-based review opt-in, excluding WIP or generated PR titles, or requesting reviews manually when the PR is ready. If your team needs uninterrupted high-volume reviews, an organization admin can enable usage-based reviews. How do review limits work?CodeRabbit enforces per-developer PR review limits for each organization. Most developers receive the normal plan review availability. For paid Pro and Pro+ PR reviews, CodeRabbit uses adaptive limits for sustained high-volume activity. When a developer's recent PR review activity reaches the 95th percentile or higher among CodeRabbit users, additional reviews become available more gradually as earlier reviews age out of the rolling window. Please refer docs for additional details. 📝 WalkthroughWalkthroughChangesSYCL model execution refactor
Estimated code review effort: 5 (Critical) | ~120 minutes Sequence Diagram(s)sequenceDiagram
participant SpatialModel
participant groupedForAll
participant LocalAllocator
participant fatalErrorFlag
SpatialModel->>groupedForAll: execute per-cell mutation callback
groupedForAll->>LocalAllocator: allocate group-local mutation storage
LocalAllocator->>fatalErrorFlag: record allocation exhaustion
groupedForAll-->>SpatialModel: complete kernel and report fatal state
🚥 Pre-merge checks | ✅ 4 | ❌ 1❌ Failed checks (1 warning)
✅ Passed checks (4 passed)
✨ Finishing Touches🧪 Generate unit tests (beta)
Thanks for using CodeRabbit! It's free for OSS, and your support helps us grow. If you like it, consider giving us a shout-out. Comment |
There was a problem hiding this comment.
Pull request overview
This PR refactors EcoLab’s SYCL integration to support a singleton-style device allocator and revised device/host allocation patterns, while updating the spatial model implementation to use new allocator types and revised kernel execution helpers.
Changes:
- Introduces new SYCL abstraction headers (
include/sycl.h,include/non-sycl.h) and a device memory allocator (include/DeviceAllocator.h) intended to support a singleton device allocator model. - Refactors
ModelData/EcolabPointallocation and several SYCL execution paths (notablygroupedForAll, mutation and generation paths) to use the new allocator approach and cached off-diagonal indexing. - Updates sample/test inputs and model parameters (shell + Python) and adjusts sparse matrix utilities accordingly.
Reviewed changes
Copilot reviewed 13 out of 13 changed files in this pull request and generated 8 comments.
Show a summary per file
| File | Description |
|---|---|
| test/00/predator-prey.sh | Updates expected trace output data for a test script. |
| src/ecolab.cc | Changes SYCL queue construction to in_order for debugging. |
| models/spatial_ecolab.py | Modifies spatial model parameters (species count, mutation, grid, timing loops). |
| models/ecolab_model.h | Refactors model/cell allocator strategy; introduces computeODiagIdx() and SYCL RNG changes. |
| models/ecolab_model.cc | Updates generation/mutation/condense logic to new data structures and device execution patterns. |
| include/usmAlloc.h | Adds conditional SYCL include to support SYCL types. |
| include/sycl.h | Adds consolidated SYCL helpers, DeviceType, allocators, RNG engine, barriers/leaders. |
| include/sparse_mat.h | Adjusts sparse matrix constructors and row-size type changes. |
| include/non-sycl.h | Adds non-SYCL fallback types for SYCL abstractions. |
| include/ecolab.h | Refactors graph parallel traversal APIs to pass (cell, idx) and revises groupedForAll scheduling. |
| include/DeviceAllocator.h | Adds a pooled device allocator + local allocator primitives. |
| include/arrays.h | Refactors array ref-counting and device-friendly assignment/copy behavior. |
💡 Add Copilot custom instructions for smarter, more guided reviews. Learn how to get started.
| // 4. Maximum number of compute units (Execution units / DSS count) | ||
| uint32_t max_compute_units = dev.get_info<sycl::info::device::max_compute_units>(); | ||
|
|
||
| size_t workGroupSize=256; |
Owner
Author
There was a problem hiding this comment.
I'm still trying to figure out how to determine best work group size. I'm thinking of making this a runtime configurable parameter. But requires many experiments.
![coderabbitai[bot]](https://avatars.githubusercontent.com/in/347564?s=60&v=4){kind=small}
There was a problem hiding this comment.
Actionable comments posted: 10
Caution
Some comments are outside the diff and can’t be posted inline due to platform limitations.
⚠️ Outside diff range comments (2)
models/ecolab_model.cc (2)
230-248: 🩺 Stability & Availability | 🟠 Major | ⚡ Quick winSkip
ModelData::mutatewhen the gathered mutation is empty.The MPI branch calls
ModelData::mutate(new_sp)before checkingcell_ids. With an extinct/empty model, this reachesmax(species)despite there being no mutation.Guard the root-side mutation with
new_sp.size()before broadcasting.🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the rest with a brief reason, keep changes minimal, and validate. In `@models/ecolab_model.cc` around lines 230 - 248, In the MPI root-side flow, guard the ModelData::mutate call with a non-empty new_sp check so empty gathered mutations are skipped. Update the block following b.gather(0) while preserving the existing resize, broadcast, and empty cell_ids return behavior.
709-718: 🗄️ Data Integrity & Integration | 🟠 Major | ⚡ Quick winResize each cell to the computed common species count.
The function computes
nspbut only updatesModelData; no cell density is padded to that size, violating the invariant stated by the comment and leaving later model arrays misaligned with cells.Proposed fix
ModelData::makeConsistent(nsp); + groupedForAll([nsp](EcolabCell& cell, size_t) { + cell.density.resize(nsp, 0); + }); syncThreads();🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the rest with a brief reason, keep changes minimal, and validate. In `@models/ecolab_model.cc` around lines 709 - 718, Update SpatialModel::makeConsistent after computing the common nsp so every cell’s density array is resized to nsp, padding newly added species with zero density. Preserve the MPI-reduced count and existing ModelData::makeConsistent and syncThreads calls, ensuring cell densities and model arrays use the same species count.
🤖 Prompt for all review comments with AI agents
Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.
Inline comments:
In `@include/arrays.h`:
- Around line 1637-1641: Update Allocator::allocator(const Allocator& alloc) to
migrate the array storage when alloc differs from m_allocator: create storage
using alloc, copy the existing elements into it, swap the new storage into the
object, and release the previous storage through its original allocator.
Preserve the fast return when alloc already matches m_allocator.
In `@include/DeviceAllocator.h`:
- Around line 41-50: Replace the unsafe head/tail logic in enqueue and dequeue
with a real bounded concurrent queue implementation, using per-slot sequence
numbers for MPMC reservation and publication or a lock that serializes both
operations. Ensure producers publish each item only after writing its slot, and
consumers reserve each item exactly once only after publication, while
preserving the empty-queue signal behavior.
- Around line 114-117: Update deviceAllocator() to call syclQ().wait_and_throw()
after the one-time DeviceAllocator::init() setup and before returning
*deviceAllocator, ensuring queued initialization kernels complete before host
allocations use the singleton.
In `@include/ecolab.h`:
- Around line 188-209: Update the fatal-error tracking in the work-group loop
around fatalErrorFlag() to use a shared atomic integer and record failures with
fetch_or(1), preventing successful groups from clearing an existing failure.
Remove the asynchronous host-side fatalError check from the loop condition, and
inspect the atomic result only after syclQ().wait_and_throw() before throwing
the existing runtime error.
- Around line 159-186: Update the launch-geometry setup before groupedForAll:
clamp workGroupSize to max_wg_size and preserve valid SIMD alignment where
applicable, then ensure wg_per_compute_unit and the resulting num_work_groups
are nonzero before launching. Handle the zero-batch case by avoiding
groupedForAll rather than constructing an nd_range with a zero global size.
In `@include/sycl.h`:
- Around line 78-83: Replace the undeclared numRngs reference in operator()’s
_OPENMP branch with the generator collection’s size, using rngs.size()
consistently with the SYCL branch while preserving the existing thread-based
indexing.
- Around line 52-57: Update SyclQAllocator’s rebind<U> in include/sycl.h lines
52-57 to alias SyclQAllocator<U, UA>, preserving the allocation strategy while
rebinding the element type. Apply the same correction to the corresponding
rebind in include/DeviceAllocator.h lines 143-148.
In `@models/ecolab_model.cc`:
- Around line 98-105: Update the iteration loop around the array_ns::map lambda
to use separate source and destination density buffers: read only from the
source buffer, write computed values to the destination buffer, synchronize the
group before exchanging them, and swap the buffers after each step. Preserve
simultaneous-update semantics and copy the final buffer back to lDensity if
needed.
- Around line 270-286: Fix the bucket construction in the LocalArray new_sp
mapping: iterate each bucket using the count offsets[i+1] - offsets[i], not
offsets[i+1], while preserving the starting position offsets[i]. Also guard the
prefix-sum finalization for nsp == 0 so it never reads speciations[nsp-1] or
offsets[-1], and retain the empty-result behavior when the total count is zero.
In `@models/ecolab_model.h`:
- Line 143: Update EcolabModel::seed so each work-item initializes the random
engine associated with its own cell index, rather than calling
SyclRandomEngine::seed(x) without an index. Use the index provided by forAll and
preserve the requested seed value for every per-cell engine.
---
Outside diff comments:
In `@models/ecolab_model.cc`:
- Around line 230-248: In the MPI root-side flow, guard the ModelData::mutate
call with a non-empty new_sp check so empty gathered mutations are skipped.
Update the block following b.gather(0) while preserving the existing resize,
broadcast, and empty cell_ids return behavior.
- Around line 709-718: Update SpatialModel::makeConsistent after computing the
common nsp so every cell’s density array is resized to nsp, padding newly added
species with zero density. Preserve the MPI-reduced count and existing
ModelData::makeConsistent and syncThreads calls, ensuring cell densities and
model arrays use the same species count.
🪄 Autofix (Beta)
Fix all unresolved CodeRabbit comments on this PR:
- Push a commit to this branch (recommended)
- Create a new PR with the fixes
ℹ️ Review info
⚙️ Run configuration
Configuration used: defaults
Review profile: CHILL
Plan: Pro
Run ID: 25c63baa-e420-495c-9346-34ba939e25e5
📒 Files selected for processing (13)
graphcodeinclude/DeviceAllocator.hinclude/arrays.hinclude/ecolab.hinclude/non-sycl.hinclude/sparse_mat.hinclude/sycl.hinclude/usmAlloc.hmodels/ecolab_model.ccmodels/ecolab_model.hmodels/spatial_ecolab.pysrc/ecolab.cctest/00/predator-prey.sh
Comment on lines
+159
to
+186
| auto dev=syclQ().get_device(); | ||
| // 1. Max threads per single work-group | ||
| size_t max_wg_size = dev.get_info<sycl::info::device::max_work_group_size>(); | ||
|
|
||
| // 2. Hardware SIMD width (Sub-group size, usually 8, 16, or 32 on Intel) | ||
| std::vector<size_t> sg_sizes = dev.get_info<sycl::info::device::sub_group_sizes>(); | ||
| size_t native_sg_size = sg_sizes.back(); // Usually 16 or 32 is best for compute | ||
|
|
||
| // 3. Total physical SLM available per hardware compute unit (DSS) | ||
| size_t max_slm_size = dev.get_info<sycl::info::device::local_mem_size>(); | ||
|
|
||
| // 4. Maximum number of compute units (Execution units / DSS count) | ||
| uint32_t max_compute_units = dev.get_info<sycl::info::device::max_compute_units>(); | ||
|
|
||
| size_t workGroupSize=256; | ||
| // if (workGroupSize > max_wg_size) | ||
| // workGroupSize = max_wg_size; // Fallback for limited devices | ||
| // | ||
| // // Ensure it aligns perfectly with hardware SIMD lanes | ||
| // workGroupSize = (workGroupSize / native_sg_size) * native_sg_size; | ||
|
|
||
| size_t wg_per_compute_unit = max_slm_size / LocalAllocatorSize; | ||
| // To maximize latency hiding, it's often beneficial to double or triple this | ||
| // so the GPU can switch to a waiting wave while another wave is blocked by a barrier. | ||
| size_t num_work_groups = max_compute_units * wg_per_compute_unit; | ||
|
|
||
| num_work_groups=std::min(num_work_groups,this->size()); | ||
| //std::cout<<max_slm_size<<" max_slm_size "<<max_compute_units<<" max_compute_units "<<num_work_groups<<" work groups of "<<workGroupSize<<" threads"<<std::endl; |
There was a problem hiding this comment.
🩺 Stability & Availability | 🟠 Major | ⚡ Quick win
🧩 Analysis chain
🏁 Script executed:
#!/bin/bash
rg -n -C4 'max_work_group_size|local_mem_size|LocalAllocatorSize|groupedForAll' include models srcRepository: highperformancecoder/ecolab
Length of output: 7087
🏁 Script executed:
#!/bin/bash
sed -n '157,210p' include/ecolab.hRepository: highperformancecoder/ecolab
Length of output: 2800
Clamp the launch geometry before groupedForAll
include/ecolab.h:159-186
workGroupSize = 256ignoresmax_wg_size, so devices with smaller limits can reject the kernel.wg_per_compute_unitcan become0whenlocal_mem_size < LocalAllocatorSize, which leaves the loop with a zero step and launches an invalidnd_range.
Restore the device-limit clamp and reject zero-sized work-group batches before launching.
🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.
In `@include/ecolab.h` around lines 159 - 186, Update the launch-geometry setup
before groupedForAll: clamp workGroupSize to max_wg_size and preserve valid SIMD
alignment where applicable, then ensure wg_per_compute_unit and the resulting
num_work_groups are nonzero before launching. Handle the zero-batch case by
avoiding groupedForAll rather than constructing an nd_range with a zero global
size.
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
This change is
Summary by CodeRabbit