Conversation
|
Auto-sync is disabled for draft pull requests in this repository. Workflows must be run manually. Contributors can view more details about this message here. |
|
/ok to test 2cc52b9 |
This comment has been minimized.
This comment has been minimized.
|
@coderabbitai review |
✅ Action performedReview finished.
|
There was a problem hiding this comment.
🔵 Needs a closer look
It introduces broad private-compiler integration and remains a draft pending exact-head CUDA 12/13 and Windows CI qualification.
Pull request overview
Adds the first executable cuda.coop backend, lowering cooperative CUB Block Load/Store through Numba-CUDA-MLIR while retaining a backend-neutral portable API.
Changes:
- Adds portable APIs, descriptors, planning contracts, typing stubs, and bundled-header resolution.
- Implements transactional Numba-CUDA-MLIR activation, rewriting, lowering, provenance tracking, and artifact caching.
- Adds packaging, documentation, examples, extensive tests, and CI integration.
File summaries
| File | Description |
|---|---|
python/cuda_coop/cuda/coop/_core/** |
Portable API and planning core |
python/cuda_coop/cuda/coop/_headers/** |
Header and toolkit resolution |
python/cuda_coop/cuda/coop/numba_mlir/** |
Compiler backend and lowering |
python/cuda_coop/cuda/coop/*.py* |
Public package exports and typing |
python/cuda_coop/tests/contracts/** |
Core and header contracts |
python/cuda_coop/tests/backends/numba_mlir/** |
Backend unit, compile, and runtime tests |
python/cuda_coop/tests/packaging/** |
Wheel, metadata, installation, and typing tests |
python/cuda_coop/tests/support/** |
Shared test utilities |
python/cuda_coop/examples/numba_mlir/** |
Block Load/Store examples |
python/cuda_coop/CMakeLists.txt |
Header bundling and provenance |
python/cuda_coop/pyproject.toml |
Package and dependency metadata |
python/cuda_coop/README.md |
Usage and compatibility guidance |
python/cuda_coop/LICENSE |
Package licensing |
ci/build_cuda_coop_python.sh |
Linux wheel build |
ci/windows/build_cuda_coop_python.ps1 |
Windows wheel build |
ci/test_cuda_coop_python.sh |
Staged backend validation |
ci/validate_cuda_coop_wheel.py |
Wheel-content validation |
ci/matrix.yaml |
CI coverage matrix |
ci/project_files_and_dependencies.yaml |
Change-impact mapping |
ci/test/inspect_changes/*cuda_coop* |
Change-detection fixtures |
ci/test/inspect_changes/{core_dirty,libcudacxx_*}.output |
Dependency expectations |
docs/python/coop.rst |
User documentation |
docs/python/coop_api.rst |
API reference |
docs/python/index.rst |
Documentation navigation |
docs/python/api_reference.rst |
API index integration |
docs/conf.py |
Documentation configuration |
Review details
- Files reviewed: 108/158 changed files
- Comments generated: 0
- Review effort level: Balanced
💡 Add a code-review agent skill or configure MCP servers for context-aware, tailored reviews. Learn more in the docs.
|
Important Review skippedWe couldn't safely recover the incremental review. No full review was started, and the last reviewed checkpoint was preserved. Retry later, or explicitly request a full review by commenting You can disable this status message by setting the Use the checkbox below for a quick retry:
No actionable comments were generated in the recent review. 🎉 ℹ️ Recent review info⚙️ Run configurationConfiguration used: Path: .coderabbit.yaml Review profile: CHILL Plan: Enterprise Run ID: 📒 Files selected for processing (10)
🚧 Files skipped from review as they are similar to previous changes (8)
Included review availability: Your plan provides up to 12 included reviews per hour; 11 remain after this review. 📝 SummarySummary by CodeRabbit
WalkthroughAdded the Changescuda.coop package and CI
Priority: ➖ Normal Merge Risk: ⚪ Minimal · up to The assessed package, compiler, validation, and test paths do not retain an actionable regression from the previously identified concerns. The change is ready to merge after normal checks. Comment |
![coderabbitai[bot]](https://avatars.githubusercontent.com/in/347564?s=60&v=4){kind=small}
There was a problem hiding this comment.
Actionable comments posted: 7
🧹 Nitpick comments (5)
python/cuda_coop/cuda/coop/numba_mlir/_compiler/_rewrite_storage.py (1)
641-641: 📐 Maintainability & Code Quality | 🔵 Trivial | 💤 Low valuesuggestion:
source_orderis reused here as the loop target after it served as the running counter at Lines 581-592. Rename the unpacked variable, for exampleuse_order, to avoid shadowing the counter.python/cuda_coop/tests/backends/numba_mlir/unit/test_nvrtc.py (1)
15-18: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick winsuggestion: Add the standard
pytestmark = [pytest.mark.backend_numba_mlir, pytest.mark.unit]assignment to the three unit-test modules. The CI stage runs the unit directory directly, so missing markers do not currently deselect these tests. The markers preserve consistent unit-test classification and marker-based selection.python/cuda_coop/cuda/coop/numba_mlir/_compiler/_rewrite.py (1)
168-168: 📐 Maintainability & Code Quality | 🔵 Trivial | 💤 Low valuesuggestion:
createdtracking is dead.
refresh_typing_contextstarts asTrue, so the|= createdat line 177 cannot change it andself._state.typingctx.refresh()always runs. Either start fromFalseto make the invocable-creation signal meaningful, or drop the flag and callrefresh()unconditionally.python/cuda_coop/cuda/coop/_core/block/load_store.py (1)
278-283: 🎯 Functional Correctness | 🔵 Trivial | ⚡ Quick winsuggestion: Normalize
items_per_threadwithnormalize_positive_int("items_per_thread", items_per_thread). The exportedmake_block_load_specandmake_block_store_specpaths currently reject validnumbers.Integralvalues such asnp.int64(4), although the shared normalizer and other integer bindings accept them. Convert the value to a plainintbefore it reachesITEMS_PER_THREADand the semantic key.python/cuda_coop/tests/backends/numba_mlir/unit/test_data_movement.py (1)
105-111: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick winsuggestion: Assert an actual rewrite match.
_run_single_phase_to_provider_boundarycan return after recognizing only theThreadDataconstructor, so the positive load tests can pass when their coop call is not inrewrite._matches. Trackbool(rewrite._matches)for each block and assert that at least one coop call matched.
ℹ️ Review info
⚙️ Run configuration
Configuration used: Path: .coderabbit.yaml
Review profile: CHILL
Plan: Enterprise
Run ID: 1e4533f7-2cf8-4484-b1eb-eba9328a7e1b
📒 Files selected for processing (158)
ci/build_cuda_coop_python.shci/matrix.yamlci/project_files_and_dependencies.yamlci/test/inspect_changes/core_dirty.outputci/test/inspect_changes/cuda_coop_python.dirty_filesci/test/inspect_changes/cuda_coop_python.outputci/test/inspect_changes/cuda_coop_python_cudax.dirty_filesci/test/inspect_changes/cuda_coop_python_cudax.outputci/test/inspect_changes/cuda_coop_python_license.dirty_filesci/test/inspect_changes/cuda_coop_python_license.outputci/test/inspect_changes/cuda_coop_python_pyproject.dirty_filesci/test/inspect_changes/cuda_coop_python_pyproject.outputci/test/inspect_changes/cuda_coop_python_pyproject_near_miss.dirty_filesci/test/inspect_changes/cuda_coop_python_pyproject_near_miss.outputci/test/inspect_changes/cuda_coop_python_windows_ci.dirty_filesci/test/inspect_changes/cuda_coop_python_windows_ci.outputci/test/inspect_changes/libcudacxx_both.outputci/test/inspect_changes/libcudacxx_public_only.outputci/test/inspect_changes/libcudacxx_thrust.outputci/test_cuda_coop_python.shci/validate_cuda_coop_wheel.pyci/windows/build_cuda_coop_python.ps1docs/conf.pydocs/python/api_reference.rstdocs/python/coop.rstdocs/python/coop_api.rstdocs/python/index.rstpython/cuda_coop/CMakeLists.txtpython/cuda_coop/LICENSEpython/cuda_coop/README.mdpython/cuda_coop/cuda/coop/__init__.pypython/cuda_coop/cuda/coop/__init__.pyipython/cuda_coop/cuda/coop/_core/__init__.pypython/cuda_coop/cuda/coop/_core/_adapter.pypython/cuda_coop/cuda/coop/_core/_algorithm.pypython/cuda_coop/cuda/coop/_core/_auto_registration.pypython/cuda_coop/cuda/coop/_core/_bindings.pypython/cuda_coop/cuda/coop/_core/_symbols.pypython/cuda_coop/cuda/coop/_core/_types.pypython/cuda_coop/cuda/coop/_core/api/__init__.pypython/cuda_coop/cuda/coop/_core/api/__init__.pyipython/cuda_coop/cuda/coop/_core/api/_dispatch.pypython/cuda_coop/cuda/coop/_core/api/_payload.pypython/cuda_coop/cuda/coop/_core/api/load_store.pypython/cuda_coop/cuda/coop/_core/api/load_store.pyipython/cuda_coop/cuda/coop/_core/api/temp_storage.pypython/cuda_coop/cuda/coop/_core/api/temp_storage.pyipython/cuda_coop/cuda/coop/_core/api/thread_data.pypython/cuda_coop/cuda/coop/_core/api/thread_data.pyipython/cuda_coop/cuda/coop/_core/api/thread_group.pypython/cuda_coop/cuda/coop/_core/api/thread_group.pyipython/cuda_coop/cuda/coop/_core/block/__init__.pypython/cuda_coop/cuda/coop/_core/block/_common.pypython/cuda_coop/cuda/coop/_core/block/load_store.pypython/cuda_coop/cuda/coop/_core/dtype_policy.pypython/cuda_coop/cuda/coop/_core/group/__init__.pypython/cuda_coop/cuda/coop/_core/group/_contracts.pypython/cuda_coop/cuda/coop/_core/group/_dispatch.pypython/cuda_coop/cuda/coop/_core/group/_model.pypython/cuda_coop/cuda/coop/_core/group/_resolution.pypython/cuda_coop/cuda/coop/_core/group/load_store.pypython/cuda_coop/cuda/coop/_core/launch.pypython/cuda_coop/cuda/coop/_core/thread_group.pypython/cuda_coop/cuda/coop/_headers/__init__.pypython/cuda_coop/cuda/coop/_headers/_identity.pypython/cuda_coop/cuda/coop/_headers/_toolkit.pypython/cuda_coop/cuda/coop/_typing.pyipython/cuda_coop/cuda/coop/numba_mlir/__init__.pypython/cuda_coop/cuda/coop/numba_mlir/__init__.pyipython/cuda_coop/cuda/coop/numba_mlir/_compiler/__init__.pypython/cuda_coop/cuda/coop/numba_mlir/_compiler/_activation.pypython/cuda_coop/cuda/coop/numba_mlir/_compiler/_artifacts.pypython/cuda_coop/cuda/coop/numba_mlir/_compiler/_caching.pypython/cuda_coop/cuda/coop/numba_mlir/_compiler/_group_load_store.pypython/cuda_coop/cuda/coop/numba_mlir/_compiler/_group_planner.pypython/cuda_coop/cuda/coop/numba_mlir/_compiler/_group_planner_support.pypython/cuda_coop/cuda/coop/numba_mlir/_compiler/_group_planning.pypython/cuda_coop/cuda/coop/numba_mlir/_compiler/_group_rewriting.pypython/cuda_coop/cuda/coop/numba_mlir/_compiler/_numba_mlir_compat.pypython/cuda_coop/cuda/coop/numba_mlir/_compiler/_nvrtc.pypython/cuda_coop/cuda/coop/numba_mlir/_compiler/_operations.pypython/cuda_coop/cuda/coop/numba_mlir/_compiler/_parameters.pypython/cuda_coop/cuda/coop/numba_mlir/_compiler/_rewrite.pypython/cuda_coop/cuda/coop/numba_mlir/_compiler/_rewrite_arguments.pypython/cuda_coop/cuda/coop/numba_mlir/_compiler/_rewrite_group_metadata.pypython/cuda_coop/cuda/coop/numba_mlir/_compiler/_rewrite_invocables.pypython/cuda_coop/cuda/coop/numba_mlir/_compiler/_rewrite_launch.pypython/cuda_coop/cuda/coop/numba_mlir/_compiler/_rewrite_load_store.pypython/cuda_coop/cuda/coop/numba_mlir/_compiler/_rewrite_payload.pypython/cuda_coop/cuda/coop/numba_mlir/_compiler/_rewrite_provenance.pypython/cuda_coop/cuda/coop/numba_mlir/_compiler/_rewrite_storage.pypython/cuda_coop/cuda/coop/numba_mlir/_compiler/_rewrite_support.pypython/cuda_coop/cuda/coop/numba_mlir/_compiler/_scalar_provenance.pypython/cuda_coop/cuda/coop/numba_mlir/_enums.pypython/cuda_coop/cuda/coop/numba_mlir/_enums.pyipython/cuda_coop/cuda/coop/numba_mlir/_group_load_store.pypython/cuda_coop/cuda/coop/numba_mlir/_group_load_store.pyipython/cuda_coop/cuda/coop/numba_mlir/_group_marker.pypython/cuda_coop/cuda/coop/numba_mlir/_lowering/__init__.pypython/cuda_coop/cuda/coop/numba_mlir/_lowering/_core.pypython/cuda_coop/cuda/coop/numba_mlir/_lowering/_load_store.pypython/cuda_coop/cuda/coop/numba_mlir/_semantic.pypython/cuda_coop/cuda/coop/numba_mlir/_temp_storage.pypython/cuda_coop/cuda/coop/numba_mlir/_temp_storage.pyipython/cuda_coop/cuda/coop/numba_mlir/_thread_data.pypython/cuda_coop/cuda/coop/numba_mlir/_thread_data.pyipython/cuda_coop/cuda/coop/numba_mlir/_thread_group.pypython/cuda_coop/cuda/coop/numba_mlir/_thread_group.pyipython/cuda_coop/cuda/coop/numba_mlir/_types.pypython/cuda_coop/cuda/coop/numba_mlir/py.typedpython/cuda_coop/cuda/coop/py.typedpython/cuda_coop/examples/numba_mlir/block_load.pypython/cuda_coop/examples/numba_mlir/block_store.pypython/cuda_coop/pyproject.tomlpython/cuda_coop/tests/__init__.pypython/cuda_coop/tests/backends/numba_mlir/__init__.pypython/cuda_coop/tests/backends/numba_mlir/compile/__init__.pypython/cuda_coop/tests/backends/numba_mlir/compile/test_block_load_store_compile.pypython/cuda_coop/tests/backends/numba_mlir/compile/test_installed_activation_compile.pypython/cuda_coop/tests/backends/numba_mlir/runtime/__init__.pypython/cuda_coop/tests/backends/numba_mlir/runtime/test_block_load_store.pypython/cuda_coop/tests/backends/numba_mlir/unit/__init__.pypython/cuda_coop/tests/backends/numba_mlir/unit/_lazy_fake_family.pypython/cuda_coop/tests/backends/numba_mlir/unit/test_caching.pypython/cuda_coop/tests/backends/numba_mlir/unit/test_data_movement.pypython/cuda_coop/tests/backends/numba_mlir/unit/test_family_registration.pypython/cuda_coop/tests/backends/numba_mlir/unit/test_group_hierarchy.pypython/cuda_coop/tests/backends/numba_mlir/unit/test_group_lowering_plan.pypython/cuda_coop/tests/backends/numba_mlir/unit/test_invocable_identity.pypython/cuda_coop/tests/backends/numba_mlir/unit/test_invocable_lifecycle.pypython/cuda_coop/tests/backends/numba_mlir/unit/test_launch_dimension_rewrites.pypython/cuda_coop/tests/backends/numba_mlir/unit/test_nvrtc.pypython/cuda_coop/tests/backends/numba_mlir/unit/test_public_foundation.pypython/cuda_coop/tests/backends/numba_mlir/unit/test_runtime_activation.pypython/cuda_coop/tests/backends/numba_mlir/unit/test_storage_rewrites.pypython/cuda_coop/tests/backends/numba_mlir/unit/test_thread_data.pypython/cuda_coop/tests/backends/numba_mlir/unit/test_thread_groups.pypython/cuda_coop/tests/contracts/__init__.pypython/cuda_coop/tests/contracts/core/test_core_bindings.pypython/cuda_coop/tests/contracts/core/test_core_block_load_store.pypython/cuda_coop/tests/contracts/core/test_core_group_load_store.pypython/cuda_coop/tests/contracts/core/test_core_group_resolution.pypython/cuda_coop/tests/contracts/core/test_core_thread_group.pypython/cuda_coop/tests/contracts/core/test_core_types.pypython/cuda_coop/tests/contracts/headers/test_identity.pypython/cuda_coop/tests/contracts/headers/test_resolution.pypython/cuda_coop/tests/contracts/headers/test_toolkit.pypython/cuda_coop/tests/packaging/test_headers.pypython/cuda_coop/tests/packaging/test_installed.pypython/cuda_coop/tests/packaging/test_metadata.pypython/cuda_coop/tests/packaging/test_typing.pypython/cuda_coop/tests/packaging/test_wheel.pypython/cuda_coop/tests/packaging/typing/invalid_consumer.pypython/cuda_coop/tests/packaging/typing/numba_consumer.pypython/cuda_coop/tests/packaging/typing/portable_consumer.pypython/cuda_coop/tests/support/__init__.pypython/cuda_coop/tests/support/group_planning.pypython/cuda_coop/tests/support/paths.py
Included review availability: Your plan provides up to 12 included reviews per hour; 11 remain after this review.
|
Addressed the current review threads in the four signed commits through
The exact-head wheel was installed outside the source tree and passed 248 |
9421804 to
58d3e31
Compare
4c93484 to
7703d64
Compare
There was a problem hiding this comment.
Actionable comments posted: 3
🧹 Nitpick comments (5)
python/cuda_coop/tests/contracts/core/test_core_group_load_store.py (1)
363-369: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick winsuggestion: restore
_GROUP_OPERATION_FAMILIESafter this test._register_group_operation_familywrites_GridOperationto the module-level registry, and this test does not undo that write. Add a fixture that snapshots and restores the registry so test-only state does not leak into the pytest process.python/cuda_coop/cuda/coop/numba_mlir/_compiler/_group_load_store.py (1)
62-71: 📐 Maintainability & Code Quality | 🔵 Trivial | 💤 Low valuesuggestion: The allowed algorithm tokens are duplicated.
_lowering/_load_store.pylines 54-61 declares the same six tokens in a localallowedset, and_resolve_algorithmthere repeats the "only 'direct'" rejection. Export one shared frozenset and reuse it in both modules so the planner and the provider cannot drift.python/cuda_coop/tests/backends/numba_mlir/unit/test_data_movement.py (1)
1109-1123: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick winimportant: This test does not distinguish its four parameter cases.
source_kindis a free variable of the Python function, andrun_frontendonly builds IR; it does not fold the comparisons. All four branches stay in the IR for every parameter value, sovalueis always a phi over the index, the element, the NumPy cast, and the backend cast. The four cases compile identical IR, and a regression in one scalar source stays hidden because the merged phi still resolves toint32.Build one kernel per source kind so each provenance path is planned alone.
Proposed fix
- def memory(source, destination): - index = cuda.threadIdx.x - if source_kind == "index": - value = index - elif source_kind == "element": - value = source[index] - elif source_kind == "numpy-cast": - value = np.int32(index + 1) - else: - value = types.int32(index + 1) - module.store(module.this_block(), destination, value) + def index_source(source, destination): + module.store(module.this_block(), destination, cuda.threadIdx.x) + + def element_source(source, destination): + module.store(module.this_block(), destination, source[cuda.threadIdx.x]) + + def numpy_cast_source(source, destination): + module.store( + module.this_block(), destination, np.int32(cuda.threadIdx.x + 1) + ) + + def compiler_cast_source(source, destination): + module.store( + module.this_block(), destination, types.int32(cuda.threadIdx.x + 1) + ) + + memory = { + "index": index_source, + "element": element_source, + "numpy-cast": numpy_cast_source, + "compiler-cast": compiler_cast_source, + }[source_kind]python/cuda_coop/tests/backends/numba_mlir/unit/test_storage_rewrites.py (1)
10-12: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick winsuggestion: Add
pytest.importorskip("numba_cuda_mlir")before the module-level backend imports. Noconftest.pyin this backend test tree provides a collection guard, so pytest raises an import error instead of skipping whennumba_cuda_mliris unavailable.python/cuda_coop/cuda/coop/numba_mlir/_compiler/_rewrite_arguments.py (1)
73-95: 🎯 Functional Correctness | 🔵 Trivial | ⚡ Quick winsuggestion: Apply the
runtime_factory_kw_prerequisitescheck to positional runtime factory kwargs. Theloadspec declaresoob_defaultas requiringnum_valid_items. The positional branch can omit a statically resolvednum_valid_itemsfrom the seen sets, while the prerequisite loop skips positional values. This allowsoob_defaultto reach the factory without its required control.
ℹ️ Review info
⚙️ Run configuration
Configuration used: Path: .coderabbit.yaml
Review profile: CHILL
Plan: Enterprise
Run ID: 7dd8757c-bc9f-4e42-93fe-ddce8db31e48
📒 Files selected for processing (78)
ci/matrix.yamlci/validate_cuda_coop_wheel.pydocs/python/coop.rstdocs/python/coop_api.rstpython/cuda_coop/README.mdpython/cuda_coop/cuda/coop/_core/_source_dump.pypython/cuda_coop/cuda/coop/_core/api/__init__.pypython/cuda_coop/cuda/coop/_core/api/_dispatch.pypython/cuda_coop/cuda/coop/_core/api/_payload.pypython/cuda_coop/cuda/coop/_core/api/temp_storage.pypython/cuda_coop/cuda/coop/_core/api/temp_storage.pyipython/cuda_coop/cuda/coop/_core/api/thread_data.pypython/cuda_coop/cuda/coop/_core/api/thread_data.pyipython/cuda_coop/cuda/coop/_core/api/thread_group.pypython/cuda_coop/cuda/coop/_core/api/thread_group.pyipython/cuda_coop/cuda/coop/_core/block/load_store.pypython/cuda_coop/cuda/coop/_core/group/_contracts.pypython/cuda_coop/cuda/coop/_core/group/_model.pypython/cuda_coop/cuda/coop/_core/group/load_store.pypython/cuda_coop/cuda/coop/_core/thread_group.pypython/cuda_coop/cuda/coop/_headers/__init__.pypython/cuda_coop/cuda/coop/_typing.pyipython/cuda_coop/cuda/coop/numba_mlir/__init__.pypython/cuda_coop/cuda/coop/numba_mlir/__init__.pyipython/cuda_coop/cuda/coop/numba_mlir/_compiler/_caching.pypython/cuda_coop/cuda/coop/numba_mlir/_compiler/_group_errors.pypython/cuda_coop/cuda/coop/numba_mlir/_compiler/_group_load_store.pypython/cuda_coop/cuda/coop/numba_mlir/_compiler/_group_planner.pypython/cuda_coop/cuda/coop/numba_mlir/_compiler/_group_planner_support.pypython/cuda_coop/cuda/coop/numba_mlir/_compiler/_group_planning.pypython/cuda_coop/cuda/coop/numba_mlir/_compiler/_group_rewriting.pypython/cuda_coop/cuda/coop/numba_mlir/_compiler/_numba_mlir_compat.pypython/cuda_coop/cuda/coop/numba_mlir/_compiler/_nvrtc.pypython/cuda_coop/cuda/coop/numba_mlir/_compiler/_rewrite.pypython/cuda_coop/cuda/coop/numba_mlir/_compiler/_rewrite_arguments.pypython/cuda_coop/cuda/coop/numba_mlir/_compiler/_rewrite_load_store.pypython/cuda_coop/cuda/coop/numba_mlir/_compiler/_rewrite_provenance.pypython/cuda_coop/cuda/coop/numba_mlir/_compiler/_rewrite_storage.pypython/cuda_coop/cuda/coop/numba_mlir/_compiler/_rewrite_support.pypython/cuda_coop/cuda/coop/numba_mlir/_compiler/_scalar_provenance.pypython/cuda_coop/cuda/coop/numba_mlir/_group_load_store.pypython/cuda_coop/cuda/coop/numba_mlir/_group_load_store.pyipython/cuda_coop/cuda/coop/numba_mlir/_lowering/_load_store.pypython/cuda_coop/cuda/coop/numba_mlir/_temp_storage.pypython/cuda_coop/cuda/coop/numba_mlir/_temp_storage.pyipython/cuda_coop/cuda/coop/numba_mlir/_thread_data.pypython/cuda_coop/cuda/coop/numba_mlir/_thread_data.pyipython/cuda_coop/cuda/coop/numba_mlir/_thread_group.pypython/cuda_coop/cuda/coop/numba_mlir/_thread_group.pyipython/cuda_coop/examples/numba_mlir/block_store.pypython/cuda_coop/tests/backends/numba_mlir/compile/test_installed_activation_compile.pypython/cuda_coop/tests/backends/numba_mlir/compile/test_scalar_store_typing.pypython/cuda_coop/tests/backends/numba_mlir/runtime/test_block_load_store.pypython/cuda_coop/tests/backends/numba_mlir/unit/test_caching.pypython/cuda_coop/tests/backends/numba_mlir/unit/test_data_movement.pypython/cuda_coop/tests/backends/numba_mlir/unit/test_family_registration.pypython/cuda_coop/tests/backends/numba_mlir/unit/test_group_hierarchy.pypython/cuda_coop/tests/backends/numba_mlir/unit/test_group_lowering_plan.pypython/cuda_coop/tests/backends/numba_mlir/unit/test_launch_dimension_rewrites.pypython/cuda_coop/tests/backends/numba_mlir/unit/test_nvrtc.pypython/cuda_coop/tests/backends/numba_mlir/unit/test_public_foundation.pypython/cuda_coop/tests/backends/numba_mlir/unit/test_runtime_activation.pypython/cuda_coop/tests/backends/numba_mlir/unit/test_storage_rewrites.pypython/cuda_coop/tests/backends/numba_mlir/unit/test_thread_data.pypython/cuda_coop/tests/backends/numba_mlir/unit/test_thread_groups.pypython/cuda_coop/tests/contracts/core/test_core_block_load_store.pypython/cuda_coop/tests/contracts/core/test_core_group_load_store.pypython/cuda_coop/tests/contracts/core/test_core_temp_storage.pypython/cuda_coop/tests/contracts/core/test_core_thread_data.pypython/cuda_coop/tests/contracts/core/test_core_thread_group.pypython/cuda_coop/tests/contracts/core/test_source_dump.pypython/cuda_coop/tests/contracts/headers/test_resolution.pypython/cuda_coop/tests/packaging/test_headers.pypython/cuda_coop/tests/packaging/test_wheel.pypython/cuda_coop/tests/packaging/typing/invalid_consumer.pypython/cuda_coop/tests/packaging/typing/numba_consumer.pypython/cuda_coop/tests/packaging/typing/portable_consumer.pypython/cuda_coop/tests/support/group_planning.py
Included review availability: Your plan provides up to 12 included reviews per hour; 11 remain after this review.
9f74170 to
44e505b
Compare
There was a problem hiding this comment.
Actionable comments posted: 3
ℹ️ Review info
⚙️ Run configuration
Configuration used: Path: .coderabbit.yaml
Review profile: CHILL
Plan: Enterprise
Run ID: a8d3b432-b41a-4777-b841-32a3861289d3
📒 Files selected for processing (39)
ci/matrix.yamldocs/python/coop.rstpython/cuda_coop/README.mdpython/cuda_coop/cuda/coop/_core/_source_dump.pypython/cuda_coop/cuda/coop/_core/api/load_store.pypython/cuda_coop/cuda/coop/_core/api/load_store.pyipython/cuda_coop/cuda/coop/_core/group/_model.pypython/cuda_coop/cuda/coop/_core/group/load_store.pypython/cuda_coop/cuda/coop/_core/thread_group.pypython/cuda_coop/cuda/coop/numba_mlir/_compiler/_descriptor_provenance.pypython/cuda_coop/cuda/coop/numba_mlir/_compiler/_group_load_store.pypython/cuda_coop/cuda/coop/numba_mlir/_compiler/_group_planner.pypython/cuda_coop/cuda/coop/numba_mlir/_compiler/_group_planning.pypython/cuda_coop/cuda/coop/numba_mlir/_compiler/_nvrtc.pypython/cuda_coop/cuda/coop/numba_mlir/_compiler/_rewrite.pypython/cuda_coop/cuda/coop/numba_mlir/_compiler/_rewrite_provenance.pypython/cuda_coop/cuda/coop/numba_mlir/_compiler/_rewrite_storage.pypython/cuda_coop/cuda/coop/numba_mlir/_compiler/_rewrite_support.pypython/cuda_coop/cuda/coop/numba_mlir/_group_load_store.pypython/cuda_coop/cuda/coop/numba_mlir/_group_load_store.pyipython/cuda_coop/cuda/coop/numba_mlir/_lowering/_core.pypython/cuda_coop/cuda/coop/numba_mlir/_temp_storage.pypython/cuda_coop/cuda/coop/numba_mlir/_types.pypython/cuda_coop/examples/numba_mlir/block_load.pypython/cuda_coop/tests/backends/numba_mlir/compile/test_block_load_store_compile.pypython/cuda_coop/tests/backends/numba_mlir/compile/test_installed_activation_compile.pypython/cuda_coop/tests/backends/numba_mlir/compile/test_storage_diagnostics.pypython/cuda_coop/tests/backends/numba_mlir/runtime/test_block_load_store.pypython/cuda_coop/tests/backends/numba_mlir/unit/test_data_movement.pypython/cuda_coop/tests/backends/numba_mlir/unit/test_group_hierarchy.pypython/cuda_coop/tests/backends/numba_mlir/unit/test_group_lowering_plan.pypython/cuda_coop/tests/backends/numba_mlir/unit/test_nvrtc.pypython/cuda_coop/tests/backends/numba_mlir/unit/test_parameter_names.pypython/cuda_coop/tests/backends/numba_mlir/unit/test_storage_rewrites.pypython/cuda_coop/tests/backends/numba_mlir/unit/test_thread_data.pypython/cuda_coop/tests/contracts/core/test_core_group_load_store.pypython/cuda_coop/tests/contracts/core/test_core_thread_group.pypython/cuda_coop/tests/contracts/core/test_source_dump.pypython/cuda_coop/tests/packaging/typing/portable_consumer.py
Included review availability: Your plan provides up to 12 included reviews per hour; 11 remain after this review.
|
For the remaining non-inline notes from the September 11 review, 886e6e0 isolates the test-only group registry and makes the four scalar provenance cases use separate kernels. I kept three suggestions out of this patch:
The six inline findings have individual replies and resolutions. Local source and installed-wheel checks pass; CI for the published head is still pending. |
This comment has been minimized.
This comment has been minimized.
f723981 to
f43a204
Compare
tpn
changed the title
This comment has been minimized.
This comment has been minimized.
This comment has been minimized.
This comment has been minimized.
This comment has been minimized.
This comment has been minimized.
Introduce the backend-neutral API, cooperative group planning, and bundled CCCL headers in a dependency-free Python package. Provide explicit backend registration without bundling a compiler adapter. Signed-off-by: Trent Nelson <trent@trent.me>
f43a204 to
bd1da34
Compare
This comment has been minimized.
This comment has been minimized.
🥳 CI Workflow Results🟩 Finished in 3h 28m: Pass: 100%/556 | Total: 7d 07h | Max: 2h 26m | Hits: 99%/855833See results here. |
Why this is needed
Python kernel compilers need shared contracts for cooperative operations:
which threads participate, how items are distributed across them, and what
temporary storage each algorithm requires. This introduces those contracts
in
cuda-coop, with the compiler adapter following separately in #11200.What changes
cuda.coopAPI, group descriptors,ThreadDataandTempStorageinterfaces, compiler-independent planning, and bundled CCCLheaders. The base
pip install cuda-coophas no Python package dependencies.direct,striped,vectorize,transpose,warp_transpose, andwarp_transpose_timesliced.16, and 32, using
direct,striped,vectorize, ortranspose.three algorithms are storage-free; transpose algorithms carry the required
allocation, alignment, and synchronization contracts.
"numba_cuda_mlir"is an accepted spelling. Registration imports the requestedadapter lazily and reports an unavailable adapter clearly. This PR contains
no compiler adapter; #11200 supplies the Numba-CUDA-MLIR implementation.
loadfills its supplied output in place and returnsNone.storealsoreturns
None. Both accept an element offset and a valid-item count relativeto the selected group's tile. Group descriptors remain usable from ordinary
Python; cooperative operations need a supported kernel compiler.
This PR temporarily targets #11523 so CI includes its Windows HostJIT build fix.
Validation
Fresh installed-wheel checks at
bd1da34d032b:Wheel validation and source, installed-package, and bundled-header comparisons passed.
Remote CI passed at
bd1da34d032b: GitHub Actions run. All exact-head checks and statuses, including pre-commit, are successful or conditionally skipped.A HostJIT histogram benchmark failed on the first attempt and passed on one targeted retry without source changes. The retry passed all 70 examples and 125 host benchmarks; the original failure remains recorded in attempt 1.