GPU dynamic memory allocation - #18
Conversation
…ethods cannot be used in an envirionment without full threads in a group participation.
using a cas loop to avoid race condition of empty stack test.
…f EcolabPoint::mutate.
Clean up device allocators. Add profiling output to spatial_ecolab
|
Warning Review limit reached
Next review available in: 27 minutes You've used all free OSS reviews for now. Wait for the free limit to reset to keep reviewing this public repository. 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. Review details⚙️ Run configurationConfiguration used: defaults Review profile: CHILL Plan: Pro Plus Run ID: 📒 Files selected for processing (2)
📝 WalkthroughWalkthroughThe change replaces the device allocation queue with active and discard stacks, adds SYCL synchronization and atomic reference handling for arrays, updates spatial mutation to use allocator-aware arrays, and changes the spatial benchmark configuration and timing output. ChangesSYCL allocation and spatial model
Classdesc revision
Estimated code review effort: 4 (Complex) | ~45 minutes Sequence Diagram(s)sequenceDiagram
participant SpatialModel
participant EcolabGraph
participant SYCLQueue
participant DeviceAllocator
SpatialModel->>EcolabGraph: submit grouped mutation work
EcolabGraph->>SYCLQueue: launch work groups
EcolabGraph->>SYCLQueue: wait for completion
EcolabGraph->>DeviceAllocator: recycle discard pile
Possibly related PRs
🚥 Pre-merge checks | ✅ 4 | ❌ 1❌ Failed checks (1 warning)
✅ Passed checks (4 passed)
✨ Finishing Touches 💡 1📝 Generate docstrings 💡
🧪 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 advances the project’s GPU/SYCL support by refactoring the device-side allocator and related array semantics, and updates the spatial model to better operate in device-parallel contexts (plus adds timing instrumentation in the spatial Python driver).
Changes:
- Reworked the SYCL device allocator to use a stack/discard-pile approach and added discard recycling after kernel execution.
- Updated
arrayref-counting / barriers and SYCL helpers (e.g.,localThreadId) to support device compilation paths. - Adjusted spatial model routines (
nsp,mutate) for grouped execution and added per-step timing output inmodels/spatial_ecolab.py.
Reviewed changes
Copilot reviewed 9 out of 9 changed files in this pull request and generated 3 comments.
Show a summary per file
| File | Description |
|---|---|
| models/spatial_ecolab.py | Changes grid sizing and adds timing + output logging for step execution. |
| models/ecolab_model.h | Updates SYCL includes and adjusts method signatures for device-aware paths. |
| models/ecolab_model.cc | Refactors nsp() and mutate() implementations for grouped/device execution. |
| include/sycl.h | Replaces onDevice() with localThreadId() helper for device code. |
| include/non-sycl.h | Adjusts non-SYCL stubs/aliases (moves LocalAllocator alias here). |
| include/ecolab.h | Adds extra SYCL synchronization and allocator discard recycling after kernels. |
| include/DeviceAllocator.h | Major allocator refactor: queue→stack, discard pile, allocator construction changes. |
| include/arrays.h | Refactors COW ref-count access to be SYCL-safe (atomic/barrier changes). |
Suppressed comments (1)
include/DeviceAllocator.h:70
Stack::popusest++and assignst=size, butsycl::atomic_refdoesn’t support these operators/assignments. This won’t compile and also can’t safely restoretopon underflow. Usefetch_add/fetch_subwith a bounds check instead.
unsigned pop()
{
Atomic t(top);
unsigned p=t++;
if (p>=size) {t=size; return ~0;} // stack empty
return slots[p];
}
💡 Add a code-review agent skill or configure MCP servers for context-aware, tailored reviews. Learn more in the docs.
| from time import perf_counter | ||
|
|
||
| epoch=2000000 | ||
| mut_factor=1000 | ||
|
|
||
| extinctions=0 | ||
| migrations=0 | ||
|
|
||
| out=open("time.dat","w") | ||
|
|
![coderabbitai[bot]](https://avatars.githubusercontent.com/in/347564?s=60&v=4){kind=small}
There was a problem hiding this comment.
Actionable comments posted: 9
🧹 Nitpick comments (3)
include/ecolab.h (1)
156-156: 📐 Maintainability & Code Quality | 🔵 TrivialTrack the missing
constoverload.This TODO leaves read-only graph traversal unavailable through
forAll. Add a const-safe overload or track this work in an issue.Do you want me to prepare the overload and its tests?
🤖 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` at line 156, Implement the missing const-safe overload for forAll using a free template function that preserves constness of the graph’s this pointer and passes it through to the functional’s arguments; add coverage for read-only graph traversal, or document the work in an issue if the overload cannot be implemented now.include/arrays.h (1)
1712-1715: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick winRemove the obsolete condition comment.
onDevice()has no remaining declarations, definitions, or callers. Usedt=x.dt;directly.🤖 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/arrays.h` around lines 1712 - 1715, In the surrounding code, remove the obsolete commented-out condition and retain the unconditional assignment `dt=x.dt;` directly before `incrRef()`.include/non-sycl.h (1)
26-33: 📐 Maintainability & Code Quality | 🔵 Trivial | 💤 Low valueDelete the dead
GroupLocalblock. The current include graph keeps theLocalAllocatoraliases separate.DeviceAllocator.his included only whenSYCL_LANGUAGE_VERSIONis defined, whilenon-sycl.his included only in the non-SYCL branch.No alias-collision fix is required for the current include graph.
🤖 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/non-sycl.h` around lines 26 - 33, Remove the commented-out GroupLocal template block immediately before SyclRandomEngine in non-sycl.h. Leave SyclRandomEngine and LocalAllocator unchanged; no alias-collision changes are needed.
🤖 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 1558-1572: Update both affected sites in
include/arrays.h:1558-1572 and include/arrays.h:1653-1668. In the array cleanup
path around ref() and decrRef(), reduce the dt-null condition across the group
before branching, move the barrier after the branch so it is unconditional, and
preserve the existing reference-count cleanup. In the second path, replace the
early return with a group-uniform condition and ensure the barrier executes
unconditionally so every work item reaches the same number of barriers.
- Around line 1595-1601: Update copy() to check the result of alloc(sz) before
dereferencing it, matching the existing null guard in resize(). If allocation
fails, preserve the allocator’s fatal-error behavior and return without
accessing dt->dt; only perform the SYCL assignment or memcpy when the allocation
succeeds.
In `@include/DeviceAllocator.h`:
- Line 200: Update the comment on LocalAllocatorSize to accurately state that
the value is 8 KiB, replacing the incorrect 32 KiB description while preserving
the constant and its occupancy-related usage.
- Around line 38-70: Update the Stack comment to describe a bounded LIFO stack,
its push/pop operations, and the constraint that concurrent device pushes and
pops must not target the same stack. In Stack::push, atomically validate that
top is greater than zero before decrementing and writing; reject or otherwise
safely handle underflow without indexing slots. Preserve the existing full-stack
behavior and avoid publishing an invalid slot.
- Around line 171-178: Update GlobalDeviceAllocator’s __SYCL_DEVICE_ONLY__
construction path so it cannot leave allocator as nullptr: delete the device
default constructor or make allocation fail in a controlled manner before
dereferencing allocator. Ensure allocate() handles any unavailable device
allocator safely while preserving the existing host initialization through
deviceAllocator().
- Around line 52-57: Update DeviceAllocator’s Stack::init to accept the
simple-range kernel’s sycl::id<1> or sycl::item<1> argument instead of calling
syclItem(). Derive the starting index and stride from that argument and the
kernel launch size, while preserving the existing slot initialization and top
reset behavior.
In `@include/sycl.h`:
- Around line 141-144: The device branch of groupLeader() must return the SYCL
group’s actual leader status instead of falling through to unconditional true;
restore an explicit device return such as syclGroup().leader() so exactly one
work item is selected. In include/sycl.h lines 141-144, update groupLeader(); in
include/DeviceAllocator.h lines 137-145, verify the discard push remains guarded
by groupLeader() and requires no direct change if corrected by that root fix.
In `@models/ecolab_model.h`:
- Line 145: Restore const access for SpatialModel::nsp by adding a
const-qualified overload while preserving the existing non-const overload and
behavior. Ensure callers with a const SpatialModel can invoke nsp without
qualification errors.
In `@models/spatial_ecolab.py`:
- Line 67: Make the time.dat handling in the timing flow rank-safe: guard the
open operation and writes around out and the lines 84-85 timing output with
myid()==0, so only rank 0 creates and writes the shared file. Preserve the
existing timing behavior while preventing other ranks from truncating or
overwriting it.
---
Nitpick comments:
In `@include/arrays.h`:
- Around line 1712-1715: In the surrounding code, remove the obsolete
commented-out condition and retain the unconditional assignment `dt=x.dt;`
directly before `incrRef()`.
In `@include/ecolab.h`:
- Line 156: Implement the missing const-safe overload for forAll using a free
template function that preserves constness of the graph’s this pointer and
passes it through to the functional’s arguments; add coverage for read-only
graph traversal, or document the work in an issue if the overload cannot be
implemented now.
In `@include/non-sycl.h`:
- Around line 26-33: Remove the commented-out GroupLocal template block
immediately before SyclRandomEngine in non-sycl.h. Leave SyclRandomEngine and
LocalAllocator unchanged; no alias-collision changes are needed.
🪄 Autofix
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 Plus
Run ID: 596c5d58-bc1f-49ef-be56-766edc52e592
📒 Files selected for processing (9)
classdescinclude/DeviceAllocator.hinclude/arrays.hinclude/ecolab.hinclude/non-sycl.hinclude/sycl.hmodels/ecolab_model.ccmodels/ecolab_model.hmodels/spatial_ecolab.py
| groupBarrier(); | ||
| if (dt) | ||
| { | ||
| if (ref()==1) | ||
| unsigned refCnt=ref(); | ||
| #ifdef __SYCL_DEVICE_ONLY__ | ||
| // ensure absolute consistency of the reference count | ||
| refCnt=sycl::reduce_over_group(syclGroup(),refCnt,sycl::minimum<unsigned>()); | ||
| #endif | ||
| if (refCnt==1) | ||
| { | ||
| free(dt); | ||
| return; | ||
| } | ||
| decrRef(); | ||
| dt=nullptr; | ||
| } else | ||
| decrRef(); | ||
| groupBarrier(); |
There was a problem hiding this comment.
🩺 Stability & Availability | 🔴 Critical | 🏗️ Heavy lift
Work-group barriers run inside data-dependent branches in array. Both sites call groupBarrier() from control flow that is not proven uniform across the work group. swap() states that an array may be per work item or shared by the group, so dt, size() and ref() can differ between work items. A barrier that only some work items reach is undefined behaviour and normally hangs the kernel. Hoist every barrier to unconditional control flow, or reduce the branch condition over the group first so all work items take the same path.
include/arrays.h#L1558-L1572: move the barrier on line 1572 out of theif (dt)block, and reduce the null test ofdtover the group before branching.include/arrays.h#L1653-L1668: replace the earlyreturnon line 1654 with a group-uniform condition, and make the barrier on line 1661 unconditional so both paths execute the same barrier count.
📍 Affects 1 file
include/arrays.h#L1558-L1572(this comment)include/arrays.h#L1653-L1668
🤖 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/arrays.h` around lines 1558 - 1572, Update both affected sites in
include/arrays.h:1558-1572 and include/arrays.h:1653-1668. In the array cleanup
path around ref() and decrRef(), reduce the dt-null condition across the group
before branching, move the barrier after the branch so it is unconditional, and
preserve the existing reference-count cleanup. In the second path, replace the
early return with a group-uniform condition and ensure the barrier executes
unconditionally so every work item reaches the same number of barriers.
There was a problem hiding this comment.
The only way for dt to be null is for the array to be created that way, for alloc to be called with zero as an argument or for line 1569 to be hit. In the first two cases, all threads participate. In the last case, the reduce_over_group is an implied barrier, so all threads must have reached that point before line 1569 is hit. In which case, all threads will reach the groupBarrier on line 1572.
There was a problem hiding this comment.
@highperformancecoder, agreed for array::release().
With the stated dt lifecycle invariant, if (dt) is work-group uniform. In the free path, all work items must participate in sycl::reduce_over_group before any work item can set dt to nullptr. Therefore, all work items reach the barrier after the branch.
I withdraw the include/arrays.h release() portion of the finding. The separate resize() path still needs its own uniformity argument.
You are interacting with an AI system.
| void init() { | ||
| top=0; // full stack | ||
| for (unsigned i=syclItem().get_global_linear_id(); i<size; | ||
| i+=syclItem().get_global_range().size()) { | ||
| slots[i].value=i; | ||
| slots[i].seq=i+1; | ||
| } | ||
| i+=syclItem().get_global_range().size()) | ||
| slots[i]=i; | ||
| } |
There was a problem hiding this comment.
🩺 Stability & Availability | 🔴 Critical | ⚡ Quick win
🧩 Analysis chain
🌐 Web query:
SYCL sycl::ext::oneapi::this_work_item::get_nd_item requirements nd_range kernel simple range
💡 Result:
The function sycl::ext::oneapi::this_work_item::get_nd_item is a SYCL extension used to retrieve a sycl::nd_item object representing the current work-item within a kernel [1][2]. Requirements and Usage: 1. Kernel Type: It is designed specifically for use within ND-range kernels [3][4][5]. It is not supported for simple range kernels (kernels launched with a simple sycl::range) [4]. 2. ND-Range Launch: The kernel must be launched with a sycl::nd_range [3]. 3. Dimensionality: The template parameter Dimensions must match the dimensionality of the currently executing kernel [3]. 4. Context: It is most commonly used within "free function" kernels (a SYCL extension) where the sycl::nd_item is not automatically passed as an argument to the function, providing a way to query work-item information like global IDs or ranges from inside the kernel body [6][4][7]. When used correctly, this function returns a sycl::nd_item instance, allowing access to standard nd_item methods such as get_global_id, get_global_linear_id, get_local_id, and get_global_range [8][1][7]. Attempting to use this function in a context that was not launched as an ND-range kernel, such as a simple range kernel, is not supported and may lead to runtime errors or undefined behavior [4][5].
Citations:
- 1: intel/llvm@cb5ef36
- 2: intel/llvm@90a6f53
- 3: https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/experimental/sycl_ext_oneapi_root_group.asciidoc
- 4: https://www.iwocl.org/wp-content/uploads/3019-James-Brodman-Intel.pdf
- 5: https://github.com/intel/llvm/blob/60f6e16aff211eb6402632d8ac938bc6730dd182/sycl/doc/extensions/proposed/sycl_ext_oneapi_work_group_memory.asciidoc
- 6: [sycl free func] The example in sycl_ext_oneapi_free_function_kernels.asciidoc can't be compiled successfully intel/llvm#20751
- 7: https://intel.github.io/llvm/design/SYCL-RTC.html
- 8: https://github.khronos.org/SYCL_Reference/iface/nd_item.html
🏁 Script executed:
#!/bin/bash
set -eu
printf '%s\n' '--- candidate files ---'
git ls-files | grep -E '(^|/)(DeviceAllocator\.h|sycl\.h)$' || true
printf '%s\n' '--- relevant symbols and call sites ---'
rg -n -C 6 'class Stack|struct Stack|void init|parallel_for|syclItem|get_nd_item' include/DeviceAllocator.h include/sycl.h 2>/dev/null || true
printf '%s\n' '--- file sizes ---'
wc -l include/DeviceAllocator.h include/sycl.h 2>/dev/null || trueRepository: highperformancecoder/ecolab
Length of output: 5694
🏁 Script executed:
#!/bin/bash
set -eu
printf '%s\n' '--- DeviceAllocator.h surrounding implementation ---'
cat -n include/DeviceAllocator.h | sed -n '1,150p'
printf '%s\n' '--- init call sites ---'
rg -n -C 5 '\.init\(\)|init\(\)' --glob '!include/DeviceAllocator.h' --glob '!include/sycl.h' .
printf '%s\n' '--- parallel_for signatures in repository ---'
rg -n -C 3 'parallel_for\s*\(' --glob '*.h' --glob '*.hpp' --glob '*.cpp' --glob '*.cc' --glob '*.cxx' .Repository: highperformancecoder/ecolab
Length of output: 15838
🌐 Web query:
SYCL 2020 queue::parallel_for range overload kernel lambda argument sycl::item sycl::id simple range
💡 Result:
In SYCL 2020, the queue::parallel_for method allows you to enqueue a basic data-parallel kernel by specifying a execution range [1][2]. When using the overload that accepts a sycl::range, the kernel lambda function can accept either a sycl::id or a sycl::item as its argument [2]. sycl::id An id represents the unique index of the current work-item within the execution range [2]. It is a simple, lightweight way to access the current thread's coordinates [2][3]. Example using sycl::id: q.submit([&](sycl::handler &h) { h.parallel_for(sycl::range<1>(size), [=](sycl::id<1> idx) { // idx provides the current global index data_accessor[idx] += 5; }); }); sycl::item An item is a more feature-rich object that encapsulates both the work-item's index and the overall execution range [2][3]. It provides methods to query information about the execution space, such as the global range, without requiring additional parameters [4][5]. Example using sycl::item: q.submit([&](sycl::handler &h) { h.parallel_for(sycl::range<1>(size), [=](sycl::item<1> itm) { auto globalId = itm.get_id; // Equivalent to sycl::id auto range = itm.get_range; // Access to the range dimensions data_accessor[globalId] += 5; }); }); Key Considerations for SYCL 2020: - Offsets: Support for passing an offset to parallel_for has been deprecated in SYCL 2020 [6]. Developers are expected to perform any necessary offset arithmetic manually within the kernel lambda [6]. - Captures: When using a lambda function as a kernel, all captures should be by copy ([=]), and the lambda must not be mutable [7]. - Dimension constraints: SYCL typically supports up to 3 dimensions. For higher-dimensional needs, developers must use a 1D range and perform manual index linearization/calculation [6].
Citations:
- 1: https://registry.khronos.org/SYCL/specs/sycl-2020/html/sycl-2020.html
- 2: https://enccs.github.io/sycl-workshop/expressing-parallelism-basic/
- 3: https://www.intel.com/content/www/us/en/developer/articles/technical/how-to-extend-cpp-applications-to-gpu-with-sycl.html
- 4: https://github.khronos.org/SYCL_Reference/iface/item.html
- 5: https://www.khronos.org/files/sycl/sycl-2020-reference-guide.pdf
- 6: https://stackoverflow.com/questions/72020782/dpc-start-the-do-loop-from-1-to-n-2-using-parallel-for-range
- 7: https://github.khronos.org/SYCL_Reference/iface/defining-kernels.html
🌐 Web query:
site:github.khronos.org/SYCL-Reference queue parallel_for range id item kernel function
💡 Result:
Relevant Khronos SYCL Reference pages:
sycl::queue::parallel_for— submits a kernel over arangeornd_range; returns an event. (github.khronos.org)sycl::handler::parallel_for— shows kernels receivingsycl::item, with examples usingrangeandid. (github.khronos.org)sycl::item— provides each work-item’s ID and range viaget_id()andget_range(). (github.khronos.org)sycl::id— represents a position within a SYCL range. (github.khronos.org)- Defining kernels — explains lambda and function-object kernel functions. (github.khronos.org)
Basic pattern:
sycl::queue q;
q.submit([&](sycl::handler& h) {
h.parallel_for<class Kernel>(
sycl::range<1>(N),
[=](sycl::id<1> i) {
// kernel body; i[0] is the work-item index
});
});With sycl::item:
[=](sycl::item<1> item) {
auto i = item.get_id()[0];
}Citations:
- 1: https://github.khronos.org/SYCL_Reference/iface/queue.html?utm_source=openai
- 2: https://github.khronos.org/SYCL_Reference/iface/command-group-handler.html?utm_source=openai
- 3: https://github.khronos.org/SYCL_Reference/iface/item.html?utm_source=openai
- 4: https://github.khronos.org/SYCL_Reference/iface/id.html?utm_source=openai
- 5: https://github.khronos.org/SYCL_Reference/iface/defining-kernels.html?utm_source=openai
Pass the simple-range index to Stack::init
syclItem() calls get_nd_item<1>(), which is not supported by the simple-range kernel at lines 112-113. Initialize the slots with the kernel’s sycl::id<1> or sycl::item<1> argument and use the launch size as the stride. Otherwise, allocator initialization can fail at runtime.
🤖 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/DeviceAllocator.h` around lines 52 - 57, Update DeviceAllocator’s
Stack::init to accept the simple-range kernel’s sycl::id<1> or sycl::item<1>
argument instead of calling syclItem(). Derive the starting index and stride
from that argument and the kernel launch size, while preserving the existing
slot initialization and top reset behavior.
| DeviceAllocator<>* allocator; | ||
|
|
||
| GlobalDeviceAllocator() = default; // note: default constructor must be called on host | ||
| #ifdef __SYCL_DEVICE_ONLY__ | ||
| GlobalDeviceAllocator(): allocator(nullptr) {} // = delete; | ||
| #else | ||
| GlobalDeviceAllocator() // note: default constructor must be called on host | ||
| {allocator=&deviceAllocator();} | ||
| #endif |
There was a problem hiding this comment.
🩺 Stability & Availability | 🟠 Major | ⚡ Quick win
🧩 Analysis chain
🏁 Script executed:
#!/bin/bash
# Find device-side default construction of arrays that use GlobalDeviceAllocator.
rg -nP -C4 '\barray\s*<[^>]*(GlobalDeviceAllocator|Allocator<)[^>]*>\s+\w+\s*(;|\{\s*\})' --type=cpp --type=cc -g '!**/build/**'
# Inspect the terminating allocator and allocate() paths for null handling.
fd -t f 'DeviceAllocator.h' | xargs -r sed -n '80,200p'Repository: highperformancecoder/ecolab
Length of output: 4390
🏁 Script executed:
#!/bin/bash
set -eu
echo '--- GlobalDeviceAllocator and array references ---'
rg -n -C3 'GlobalDeviceAllocator|m_allocator|array\s*\(' include src test examples --glob '!**/build/**' 2>/dev/null || true
echo '--- arrays.h structure and constructors ---'
fd -t f 'arrays.h' . | while read -r f; do
echo "FILE: $f"
ast-grep outline "$f" | head -120
rg -n -C8 'm_allocator|array\s*\(|allocate\s*\(' "$f"
done
echo '--- array declarations and allocator aliases ---'
rg -n -C3 'using .*Allocator|typedef .*Allocator|array\s*<' . \
--glob '*.{h,hpp,c,cpp,cc,cxx}' \
--glob '!**/build/**' \
--glob '!**/dist/**' 2>/dev/null || trueRepository: highperformancecoder/ecolab
Length of output: 50384
🏁 Script executed:
#!/bin/bash
set -eu
echo '--- All direct GlobalDeviceAllocator uses ---'
rg -n -C5 'GlobalDeviceAllocator' . --glob '!**/build/**' --glob '!**/dist/**'
echo '--- SYCL allocator aliases and device array declarations ---'
fd -t f -e h -e hpp -e cpp -e cc -e cxx . include src test models | while read -r f; do
if rg -q 'DeviceAllocator|GlobalDeviceAllocator|__SYCL_DEVICE_ONLY__|array_ns::array|array<' "$f"; then
echo "FILE: $f"
rg -n -C4 'DeviceAllocator|GlobalDeviceAllocator|__SYCL_DEVICE_ONLY__|array_ns::array|array<' "$f" | head -160
fi
done
echo '--- sycl.h relevant definitions ---'
fd -t f 'sycl.h' . | while read -r f; do
echo "FILE: $f"
sed -n '1,180p' "$f"
done
echo '--- arrays.h include and allocator-selection context ---'
sed -n '1,80p' include/arrays.h
sed -n '1240,1300p' include/arrays.hRepository: highperformancecoder/ecolab
Length of output: 50383
🏁 Script executed:
#!/bin/bash
set -eu
echo '--- ModelData and EcolabPoint definitions ---'
sed -n '1,125p' models/ecolab_model.h
echo '--- DeviceType definition and uses ---'
rg -n -C8 'template.*DeviceType|struct DeviceType|class DeviceType|DeviceType<' include models src \
--glob '*.{h,hpp,c,cpp,cc,cxx}' | head -240
echo '--- Device-side default construction of GlobalDeviceAllocator-backed objects ---'
rg -n -C8 'UnsignedArray|Allocator<|cell_ids|density|array<[^;]*(GlobalDeviceAllocator|ModelData::Allocator)' \
models/ecolab_model.h models/ecolab_model.cc include \
--glob '*.{h,hpp,c,cpp,cc,cxx}' | head -320
echo '--- Read-only semantic probe of the relevant preprocessor branches ---'
python3 - <<'PY'
from pathlib import Path
allocator = Path("include/DeviceAllocator.h").read_text()
arrays = Path("include/arrays.h").read_text()
model = Path("models/ecolab_model.h").read_text()
checks = {
"device constructor initializes allocator to nullptr":
"GlobalDeviceAllocator(): allocator(nullptr) {}" in allocator,
"device allocate dereferences allocator without a guard":
"{return reinterpret_cast<T*>(allocator->allocate(n*sizeof(T)));}" in allocator,
"array default constructor is defaulted":
"array()=default;" in arrays,
"array stores allocator as a value member":
"A m_allocator;" in arrays,
"EcolabPoint uses GlobalDeviceAllocator under SYCL":
"template <class T> using Allocator=GlobalDeviceAllocator<T>;" in model,
"EcolabPoint has a GlobalDeviceAllocator-backed array member":
"array<int,Allocator<int>> density;" in model,
}
for name, result in checks.items():
print(f"{'PASS' if result else 'FAIL'}: {name}")
PYRepository: highperformancecoder/ecolab
Length of output: 32452
Reject null device allocator construction.
When GlobalDeviceAllocator<T> is default-constructed on the device, allocator is nullptr, but allocate() dereferences it without a check. A default-constructed array<T, GlobalDeviceAllocator<T>> therefore fails on its first non-zero allocation. Delete the device constructor or report a controlled allocation failure.
🤖 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/DeviceAllocator.h` around lines 171 - 178, Update
GlobalDeviceAllocator’s __SYCL_DEVICE_ONLY__ construction path so it cannot
leave allocator as nullptr: delete the device default constructor or make
allocation fail in a controlled manner before dereferencing allocator. Ensure
allocate() handles any unavailable device allocator safely while preserving the
existing host initialization through deviceAllocator().
| }; | ||
|
|
||
| constexpr static unsigned LocalAllocatorSize=30*1024; // 32KiB = half typical local storage | ||
| constexpr static unsigned LocalAllocatorSize=8*1024; // 32KiB = half typical local storage |
There was a problem hiding this comment.
📐 Maintainability & Code Quality | 🟡 Minor | ⚡ Quick win
The comment contradicts the value.
LocalAllocatorSize is now 8 KiB, but the comment states 32 KiB. include/ecolab.h derives wg_per_compute_unit from this constant, so a wrong comment misleads occupancy tuning.
📝 Proposed comment fix
- constexpr static unsigned LocalAllocatorSize=8*1024; // 32KiB = half typical local storage
+ constexpr static unsigned LocalAllocatorSize=8*1024; // 8KiB per work group of local storage📝 Committable suggestion
‼️ IMPORTANT
Carefully review the code before committing. Ensure that it accurately replaces the highlighted code, contains no missing lines, and has no issues with indentation. Thoroughly test & benchmark the code to ensure it meets the requirements.
| constexpr static unsigned LocalAllocatorSize=8*1024; // 32KiB = half typical local storage | |
| constexpr static unsigned LocalAllocatorSize=8*1024; // 8KiB per work group of local storage |
🤖 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/DeviceAllocator.h` at line 200, Update the comment on
LocalAllocatorSize to accurately state that the value is 8 KiB, replacing the
incorrect 32 KiB description while preserving the constant and its
occupancy-related usage.
| return *objects[makeId(x,y)]; | ||
| } | ||
| array<unsigned> nsp() const; | ||
| array<unsigned> nsp(); |
There was a problem hiding this comment.
🎯 Functional Correctness | 🟡 Minor | ⚡ Quick win
🧩 Analysis chain
🏁 Script executed:
#!/bin/bash
set -euo pipefail
# Map nsp declarations and definitions before inspecting call sites.
ast-grep outline models --items all --type function --match nsp
# Inspect C++ call sites and nearby const-qualified model references.
rg -n --type=cpp -C4 '(?:\.|->)nsp\s*\(|const\s+SpatialModel' .Repository: highperformancecoder/ecolab
Length of output: 710
🏁 Script executed:
#!/bin/bash
set -euo pipefail
printf '%s\n' '--- declarations and definitions ---'
rg -n -C3 '\bnsp\s*\(' models --glob '*.{h,hpp,hh,cc,cpp,cxx}' --glob '!*.md'
printf '%s\n' '--- all nsp call sites ---'
rg -n -C3 '\bnsp\s*\(' . --glob '*.{h,hpp,hh,cc,cpp,cxx,c}'
printf '%s\n' '--- SpatialModel type uses ---'
rg -n -C3 '\b(const\s+)?SpatialModel\b' . --glob '*.{h,hpp,hh,cc,cpp,cxx}'Repository: highperformancecoder/ecolab
Length of output: 7604
Preserve const on SpatialModel::nsp. No in-repository callers use this method, but removing const breaks external callers with a const SpatialModel. Add a const overload.
🤖 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.h` at line 145, Restore const access for
SpatialModel::nsp by adding a const-qualified overload while preserving the
existing non-const overload and behavior. Ensure callers with a const
SpatialModel can invoke nsp without qualification errors.
This change is
Summary by CodeRabbit
Performance and Reliability
API Improvements
Simulation Updates