Skip to content

feat(rocke): switch conv kernels to AOT - #12783

Merged
bartekxk merged 7 commits into
developfrom
users/barkocot/conv-aot
Oct 5, 2026
Merged

bartekxk merged 7 commits into
developfrom
users/barkocot/conv-aot

Conversation

@bartekxk

Copy link
Copy Markdown
Contributor

Motivation

Make conv kernels more generic.

Technical Details

Switch from JIT compilation to AOT

Test Plan

run_all.py

Test Result

pass

Submission Checklist

JIRA ID: AICK-2366

@therock-pr-bot

therock-pr-bot Bot commented Sep 30, 2026 •

Copy link
Copy Markdown

✅ All Checks Passed — Ready for Review

Check Status Details
📝 PR Description ✅ Pass —
⛔ Forbidden Files ✅ Pass —
🧪 Unit Test ✅ Pass —
🔎 pre-commit ✅ Pass —
🚫 Draft PR 🔜 To Be Enabled —
🚩 Feature Flag 🔜 To Be Enabled —
📊 Code Coverage 🔜 To Be Enabled —
🤖 therock-pr-bot ✅ Pass —

🎉 All checks passed! This PR is ready for review.

📖 Need help? See the Policy FAQ for details on every check and how to fix failures.

🙋 Wish to Override Policy?

Copilot AI left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Copilot review overview

🟡 Changes recommended

Unresolved ABI construction, manifest packing, and split-K pipeline defects can prevent launches or produce incorrect results.

Review effort: Balanced
Findings: 5 High severity · 1 Medium severity

Open (6)
What changed in this PR

Migrates rocKE convolution kernels from shape-specialized JIT compilation to shape-generic AOT kernels using runtime kernargs.

Changes:

  • Adds dynamic transforms, descriptors, pipelines, and mirrored C++ emitters.
  • Centralizes convolution ABI, manifests, caching, dispatch, and launch packing.
  • Updates convolution correctness and cache tests.
File Description
platform/​python/​rocke/​portable_ir/​drivers/​roll_nd_coverage.py Removes obsolete JIT sampling probe.
platform/​python/​rocke/​helpers/​transforms.py Adds dynamic transforms and descriptors.
platform/​python/​rocke/​helpers/​pipeline.py Adds runtime-extent ping-pong pipeline.
platform/​python/​rocke/​helpers/​manifest.py Emits AOT convolution signatures.
platform/​python/​rocke/​helpers/​conv_abi.py Defines centralized convolution ABI.
platform/​python/​rocke/​examples/​gfx950/​conv_dgrad/​run_one_dgrad.py Uses AOT dgrad arguments.
platform/​cpp/​instances/​common/​conv_implicit_gemm_dgrad.cpp Converts dgrad emission to AOT.
platform/​cpp/​instances/​common/​conv_implicit_gemm_conv_epilogue.cpp Uses runtime epilogue descriptors and bounds.
platform/​cpp/​instances/​common/​conv_implicit_gemm_conv_build_glue.cpp Declares and wires runtime parameters.
platform/​cpp/​instances/​common/​conv_direct_grouped_spec_and_validate.cpp Adds direct-convolution validation.
platform/​cpp/​instances/​common/​conv_direct_grouped_common.cpp Adds shared direct AOT plumbing.
platform/​cpp/​instances/​common/​conv_direct_grouped_build_dgrad.cpp Converts direct dgrad to AOT.
platform/​cpp/​instances/​common/​conv_direct_grouped_build_depthwise.cpp Converts depthwise forward to AOT.
platform/​cpp/​instances/​common/​conv_direct_grouped_build_depthwise_spatial.cpp Converts spatial depthwise loops to AOT.
platform/​cpp/​instances/​common/​conv_direct_grouped_build_4c_phases_and_loop.cpp Adds runtime 4c row loops.
platform/​cpp/​instances/​common/​conv_direct_grouped_build_16c_phases.cpp Adds runtime 16c descriptors.
platform/​cpp/​include/​rocke/​instance_conv_implicit_gemm.h Declares dynamic descriptor APIs.
platform/​cpp/​include/​rocke/​instance_conv_implicit_gemm_internal.h Extends build context with AOT values.
platform/​cpp/​include/​rocke/​instance_conv_implicit_gemm_dgrad.h Adds dgrad search capability bound.
platform/​cpp/​include/​rocke/​instance_conv_direct_grouped_internal.h Defines direct AOT parameter context.
platform/​cpp/​include/​rocke/​helper_rocke.helpers.transforms.h Declares dynamic transform infrastructure.
library/​tests/​test_direct_kernel_cache.py Tests direct AOT cache behavior.
library/​tests/​test_direct_conv_correctness.py Updates direct launch correctness tests.
library/​tests/​test_conv_wgrad_correctness.py Updates wgrad AOT launches.
library/​tests/​test_conv_fwd_correctness.py Updates forward AOT launches.
library/​tests/​test_conv_dgrad_correctness.py Updates dgrad AOT launches.
library/​kernels/​common/​manifest_runner/​conv.py Packs runtime convolution arguments.
library/​kernels/​common/​conv_implicit_gemm_wgrad_two_stage.py Adds two-stage AOT launch helpers.
library/​dispatch/​grouped_convolution.py Returns full AOT signatures and values.
library/​benchmarks/​common/​kernel_cache.py Adds AOT kernel identity and cache.
library/​benchmarks/​common/​grouped_conv/​benchmark_grouped_conv.py Updates grouped benchmarks for AOT.

💡 Add a code-review agent skill or configure MCP servers for context-aware, tailored reviews. Learn more in the docs.

@bartekxk

Copy link
Copy Markdown
Contributor Author

Review: #12783 — feat(rocke): switch conv kernels to AOT

Reviewed commit: ac0b2f517c3

I found five bugs: two break things outright, and three produce wrong results or wrong launches. I verified each one by reading the code, without running anything on a GPU. For the first one, pyflakes on the changed files also reports the undefined name.

1. Grouped-conv benchmark crashes with NameError

Location: library/benchmarks/common/grouped_conv/benchmark_grouped_conv.py:1001 (also lines 1261–1271 and 1318–1328)

_run_fwd and _run_wgrad read spec.tile_m, spec.tile_n, spec.tile_k and spec.split_k, but neither function has a variable named spec. The loop variables are dspec / instance_spec, so every fwd and wgrad candidate raises NameError.

Fix: use the correct loop variable.

2. Wgrad unroll_k loop runs the MFMA on the wrong buffer

Location: library/kernels/common/conv_implicit_gemm_wgrad.py:2329

Each iteration loads the next tile into the second buffer, but the MFMA always reads the first buffer, which only ever holds the first tile. With more than one K tile, dW ends up as the first tile's contribution multiplied by the tile count. The previous version alternated the buffers correctly.

Fix: switch the compute buffer on each iteration, as the previous version did.

3. Split-K wgrad can add a tile to dW twice

Location: platform/python/rocke/helpers/pipeline.py:382 (run_ping_pong_dynamic)

Each step of the loop computes a second tile and assumes anything past the loop bound reads as zero. That is only true at the real end of the tensor. The async wgrad path passes the end of its split-K slice as the bound. With an odd number of tiles per slice, it therefore also computes the first tile of the next slice.

Example: 5 K tiles split 2 ways gives 3 tiles per slice. Slice 0 also computes tile 3, which slice 1 computes too, so tile 3 is added to dW twice.

Fix: guard the second (tail) tile against the slice bound, not only the tensor bound.

4. Dispatcher launch values and signature disagree with the kernel and grid

Location: library/dispatch/grouped_convolution.py:1747 (also :1700)

  • launch_values_for and _signature_for read the raw spec.split_k and getattr(spec, "two_stage", False).
  • ConvGroupedSpec has no two_stage field, so that value is always False.
  • The kernel and the grid instead use _resolve_wgrad_split_k, which resolves auto split-K, clamps it, and can switch the kernel to two-stage.

Consequences:

  • Two-stage kernel: the signature leaves out ws_ptr/ws_bytes, so ks/ks_count get packed into the workspace-pointer slot.
  • Auto split-K resolving to, e.g., 4: the grid launches groups × 4 blocks in z, but the kernel is told there is 1 slice, so it decodes group and slice wrongly.

Fix: compute launch values and the signature from _resolve_wgrad_split_k, the same helper the kernel and grid use.

5. Grouped dgrad from the kernel cache only computes group 0

Location: library/benchmarks/common/benchmark_implicit_gemm_conv.py:3689

The launch grid is (blocks, 1, split_k), but the dgrad kernel reads the conv group from blockIdx.y. The dispatcher correctly puts the group count in grid y; this path does not. The cache sweep does build grouped dgrad kernels, so this path is reachable, and the benchmark still reports timings for the incomplete result.

Fix: launch with (blocks, groups, split_k).

Note on byte-identity

Bugs 2 and 3 are in kernel emission. The fixes have to be mirrored in the C++ engine and the golden re-blessed in the same change. Then run tools/check_byte_identity.py for both llvm20 and llvm22.

@therock-pr-bot

therock-pr-bot Bot commented Sep 30, 2026 •

Copy link
Copy Markdown

🎉 All checks passed! This PR is ready for review.

Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
@bartekxk
bartekxk force-pushed the users/barkocot/conv-aot branch from 1419d3a to f4cb7e1 Compare October 1, 2026 11:52
@bartekxk
bartekxk requested a balanced review from Copilot October 1, 2026 11:53

Copilot AI left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

@aledudek aledudek left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Could you expand the test plan? This changes the ABI and IR of every conv kernel, could you please list:

  • check_byte_identity.py at llvm20 and llvm22
  • a run with ROCKE_CPP_STRICT=1 (or the engine build_id), so we know the C++ engine actually ran
  • the on-device conv correctness tests (fwd/wgrad/dgrad/direct) and which GPU they ran on, with platform/tests and library/tests run separately

Copilot AI left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Copilot review overview

🟡 Changes recommended

The dynamic magic transform lacks full Python/C++ API parity, and several documented launch workflows are currently invalid.

Review effort: Balanced
Findings: 1 Medium severity · 1 Low severity

Open (2)
Resolved since last review (3)

@aledudek aledudek left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Looks good besides the group merge problem, left a few minor comments, good to merge after resolving that

@github-actions github-actions Bot added the rocke-lib: attention Touches rocKE library attention/FMHA kernels label Oct 5, 2026
@bartekxk
bartekxk enabled auto-merge (squash) October 5, 2026 12:45
@bartekxk
bartekxk merged commit 3c7094a into develop Oct 5, 2026
96 checks passed
@bartekxk
bartekxk deleted the users/barkocot/conv-aot branch October 5, 2026 15:03
jakpiase added a commit that referenced this pull request Oct 6, 2026
…ol-cpp-port

Adapts the column-streamed depthwise kernel (DirectDepthwiseColSpec) to the
AOT conv ABI that develop switched the direct-conv kernels to (#12783).

The dwcol kernel unrolled its input rows over the whole image, keeping a
Ho x block_w accumulator band, so its trip count baked the image height --
exactly what AOT forbids. Following the DirectConvSpec precedent, the band
is now a build-time row tile:

- new block_h knob (default 16): each block covers block_h output rows;
  block_id_z encodes n * ceil(Ho / block_h) + h_tile, with the tile count
  derived from the p_Ho kernarg; stores are bounded against p_Ho.
- the kernel declares the direct-conv kernarg block (emit_direct_params /
  rocke_dconv_emit_params) and uses the dynamic A/D descriptors, so N, H, W,
  groups and the tensor strides are runtime.
- live_f32 is now block_h * block_w + KH; n_iters is
  (block_h - 1) * stride + KH.
- ch_tile_exact / w_tile_exact guard elision is removed: the group count
  and Wo are kernargs, so the guards are always emitted.
- direct_launch_geometry handles DirectDepthwiseColSpec.

Both engines are mirrored; byte-identity is GREEN for every family at
llvm20 and llvm22. The dwcol representative-IR goldens are re-blessed
(only the dwcol entries moved).

AOT cache: direct_depthwise_col is a new forward variant in
DIRECT_CAPABILITIES. It is padding-generic (rows are bounded by p_Ho), so it
takes every PAD in [0, KH-1] except stride-1 over-padding. The branch's
per-direction depthwise block_w grids (DW_BLOCK_W_FWD / DW_BLOCK_W_DGRAD)
and the col block_h/block_w grids move into direct_kernel_sweep, which
develop made the single source for the JIT sweep and the cache.

Conflict resolution: the preload depthwise builders take develop's
loop-only form, which supersedes this branch's unroll threshold of 0.

Tests: the dwcol GPU correctness suite gains row-tile shapes and a test
that runs one compiled binary across several image shapes; the shape-
invariance test covers dwcol; the cache tests cover the new variant; the
preload tap-ceiling boundary test is re-pinned, since forward depthwise now
only accepts odd filters with "same" padding.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
jakpiase added a commit that referenced this pull request Oct 6, 2026
…ke-conv-bwd-merged-groups-heuristic

Adapts the depthwise wgrad merged-groups heuristic to the AOT switch (#12783).

Conflict was in test_grouped_conv_wgrad_dispatch.py, where both sides appended
a test class; both are kept. AOT's TestLaunchContractMatchesKernel asserted the
pre-merge invariant gridDim.z == groups * split_k, which a merged launch breaks
(z is groups/Gm * split_k). It now asks WgradConvSpec.grid_groups, the same
property _wgrad_grid divides by, so it reduces to the old assertion at Gm == 1.

grouped_convolution.py merged cleanly and is correct as merged: both AOT entry
points go through _resolve_wgrad_split_k, so a merged split_k > 1 resolves to
two-stage and the ws_ptr/ws_bytes pair joins the kernarg ABI on its own. The
launch values stay built from the true (unmerged) problem, which the kernel
expects -- it scales by Gm in-kernel -- and p_num_pid_m/n still match the merged
grid because the merge gate admits only single-tile merges. That invariant is
now documented and pinned by a test rather than left implicit.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
jakpiase added a commit that referenced this pull request Oct 6, 2026
…ke-conv-fwd-merged-groups

Brings in the switch of conv kernels from JIT to AOT compilation (#12783)
and adapts the depthwise forward group-merge feature to it.

Under AOT the problem shape is passed as kernel arguments, so a binary is
shared by every compatible shape. Following the wgrad merge precedent, the
kernargs keep describing the TRUE problem and the merged geometry is folded
in-kernel from the build-time degree: merge is depthwise-only, so the merged
cpg/kpg are exactly Gm (constants), and only the merged reduction K_gemm*Gm
is a runtime product.

- Builder (Python and C++, byte-identical): merged k decode in the split
  address form (shift/mask for the merged group, magic division for X),
  merged B load with the diagonal mask, merged K-loop extent for the plain,
  unroll_k and async drivers, merged group base and N bound in the epilogue.
  The grouped decode stays engaged at Gm == groups because the binary serves
  every group count Gm divides. 3-D is now rejected by the merge gate.
- Default load/store widths come from the merged problem in both the
  builder and the validator (_sync_load_vecs, the cshuffle default).
- Python's open-width cshuffle store now uses the kpg-derived width, as the
  C++ engine and is_valid_spec already did; the fixed 8 dropped every output
  when kpg was not a multiple of 8 (e.g. unmerged depthwise).
- AOT cache: the sweep builds merged depthwise fwd kernels (cshuffle, plain
  and unroll_k); the cache serves them to depthwise problems Gm divides with
  Gm <= tile_n; run-from-cache launches groups/Gm along z.
- Parity configs renumbered after develop's new ones (merged now 20-25, plus
  a merged unroll_k case); golden IR hashes re-blessed for the merged cases.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
@assistant-librarian

Copy link
Copy Markdown
Contributor

TheRock Submodule Bump Activity

Newest first

bartekxk pushed a commit that referenced this pull request Oct 13, 2026
## Motivation

After #12783 (JIT → AOT), many conv shapes got slower. In some cases the
AOT cache had no kernel at all. The causes:

1. The cache left out kernels the JIT sweep used to pick.
2. The shape-generic kernels compute index math with runtime divides
inside the K loop.

## Technical Details

**Cache coverage**
- **Direct-conv wgrad** is now cached.
`DirectConvWgradSpec.runtime_channels` makes one binary serve any
channel count.
- **Direct MFMA dgrad for dense problems** is now cached, through
runtime-channel modes of `DirectConvSpec` and its transpose/reorganize
helpers. The C++ engine gets ports of these.
- **Implicit-GEMM wgrad and dgrad** entries now use the K-outer LDS
layout, the same as the JIT sweep picks.
- **Non-split wgrad** (split_k=1) is cached, and split-K degrees 256/512
are tried at launch.
- **13x13 and 31x31 depthwise forward** filters are cached.

**New kernel loop options** (off by default, so existing IR is
unchanged)
- `WgradConvSpec.prefetch_addr`: computes the next tile's addresses
while the current loads are in flight.
- `DgradConvSpec.uniform_k_decode`: moves the per-element divides out of
the K loop.
- Both are turned on by deterministic, per-geometry selection rules.

**Benchmark tooling**
- Fixed a device-memory leak in `--run-from-cache` (two-stage workspaces
and dgrad sub-GEMM tables).
- `--split-k-prune` now works on the cache path.
- Early stop drops hopeless kernels after a single launch.
- The compare script no longer skips direct conv for wgrad.

The follow-up commit addresses self-review findings: C++ bounds and
validation, dgrad cache coverage, early-stop robustness, and extra
tests.

Cache size grows about 16% per dtype. Most of it is the non-split wgrad
entries.

## Test Plan

- `check_byte_identity.py`: llvm20 and llvm22. New parity configs cover
each new mode.
- `platform/tests/run_all.py` with ctest.
- CPU: `test_conv_abi`, `test_direct_kernel_cache`, `test_early_stop`,
`test_conv_wgrad_prefetch_addr`, `test_conv_dgrad_uniform_k_decode`,
layering/signature tests.
- GPU (gfx950): `test_conv_wgrad_correctness`,
`test_conv_dgrad_correctness`, `test_direct_conv_correctness`,
`test_conv_fwd_correctness`.
- `--run-from-cache --verify` on sampled wgrad, dgrad and depthwise
shapes.

## Test Result

- Byte identity is GREEN for all 73 families on both LLVM flavors.
- ctest passes 14/14.
- All the listed tests pass.
- Every new kernel mode passes `--verify` against the torch reference.
- `run_all.py` reports two failures that also occur on the base commit,
unrelated to this change: the relative-path guard on
`platform/CMakeLists.txt`, and `test_parity_matrix` for `tf32` on
gfx942.

**Performance** (relative only).

How it was measured:
- **Setup:** gfx950, NHWC, bf16 and fp16. Shapes are 297 forward, wgrad
and dgrad problems from three model shape lists.
- **Baseline:** best JIT result before #12783.
- **New side:** best kernel from this branch's AOT cache, measured with
`--run-from-cache` at 3 warmup / 10 iterations.
- **Correctness:** the new kernel modes passed `--verify` against a
reference. The other cached kernels were not re-verified in the timing
runs.

Results:
- wgrad shapes went from typically 2–3× slower than the JIT baseline
(worst case about 6×) to parity on most measured configurations.
- dgrad shapes that were 1.1–1.3× slower are about level with the
baseline.
- The shape-generic kernel loops alone got 1.05–1.3× faster on the
measured configurations.

Caveats:
- At this iteration count some rows varied by up to 1.6× between runs.
- A few remaining gaps are outside this PR:
- Direct depthwise kernels are slower in the JIT path too, which is
being addressed separately.
- dgrad with stride > 2, and direct dgrad with an odd number of
16-channel dY tiles (80, 112, …), are not yet cached.

JIRA-ID : AICK-2519

---------

Co-authored-by: Claude Opus 5.5 <noreply@anthropic.com>
mpanoop pushed a commit to mpanoop/rocm-libraries that referenced this pull request Oct 14, 2026
## Motivation

After ROCm#12783 (JIT → AOT), many conv shapes got slower. In some cases the
AOT cache had no kernel at all. The causes:

1. The cache left out kernels the JIT sweep used to pick.
2. The shape-generic kernels compute index math with runtime divides
inside the K loop.

## Technical Details

**Cache coverage**
- **Direct-conv wgrad** is now cached.
`DirectConvWgradSpec.runtime_channels` makes one binary serve any
channel count.
- **Direct MFMA dgrad for dense problems** is now cached, through
runtime-channel modes of `DirectConvSpec` and its transpose/reorganize
helpers. The C++ engine gets ports of these.
- **Implicit-GEMM wgrad and dgrad** entries now use the K-outer LDS
layout, the same as the JIT sweep picks.
- **Non-split wgrad** (split_k=1) is cached, and split-K degrees 256/512
are tried at launch.
- **13x13 and 31x31 depthwise forward** filters are cached.

**New kernel loop options** (off by default, so existing IR is
unchanged)
- `WgradConvSpec.prefetch_addr`: computes the next tile's addresses
while the current loads are in flight.
- `DgradConvSpec.uniform_k_decode`: moves the per-element divides out of
the K loop.
- Both are turned on by deterministic, per-geometry selection rules.

**Benchmark tooling**
- Fixed a device-memory leak in `--run-from-cache` (two-stage workspaces
and dgrad sub-GEMM tables).
- `--split-k-prune` now works on the cache path.
- Early stop drops hopeless kernels after a single launch.
- The compare script no longer skips direct conv for wgrad.

The follow-up commit addresses self-review findings: C++ bounds and
validation, dgrad cache coverage, early-stop robustness, and extra
tests.

Cache size grows about 16% per dtype. Most of it is the non-split wgrad
entries.

## Test Plan

- `check_byte_identity.py`: llvm20 and llvm22. New parity configs cover
each new mode.
- `platform/tests/run_all.py` with ctest.
- CPU: `test_conv_abi`, `test_direct_kernel_cache`, `test_early_stop`,
`test_conv_wgrad_prefetch_addr`, `test_conv_dgrad_uniform_k_decode`,
layering/signature tests.
- GPU (gfx950): `test_conv_wgrad_correctness`,
`test_conv_dgrad_correctness`, `test_direct_conv_correctness`,
`test_conv_fwd_correctness`.
- `--run-from-cache --verify` on sampled wgrad, dgrad and depthwise
shapes.

## Test Result

- Byte identity is GREEN for all 73 families on both LLVM flavors.
- ctest passes 14/14.
- All the listed tests pass.
- Every new kernel mode passes `--verify` against the torch reference.
- `run_all.py` reports two failures that also occur on the base commit,
unrelated to this change: the relative-path guard on
`platform/CMakeLists.txt`, and `test_parity_matrix` for `tf32` on
gfx942.

**Performance** (relative only).

How it was measured:
- **Setup:** gfx950, NHWC, bf16 and fp16. Shapes are 297 forward, wgrad
and dgrad problems from three model shape lists.
- **Baseline:** best JIT result before ROCm#12783.
- **New side:** best kernel from this branch's AOT cache, measured with
`--run-from-cache` at 3 warmup / 10 iterations.
- **Correctness:** the new kernel modes passed `--verify` against a
reference. The other cached kernels were not re-verified in the timing
runs.

Results:
- wgrad shapes went from typically 2–3× slower than the JIT baseline
(worst case about 6×) to parity on most measured configurations.
- dgrad shapes that were 1.1–1.3× slower are about level with the
baseline.
- The shape-generic kernel loops alone got 1.05–1.3× faster on the
measured configurations.

Caveats:
- At this iteration count some rows varied by up to 1.6× between runs.
- A few remaining gaps are outside this PR:
- Direct depthwise kernels are slower in the JIT path too, which is
being addressed separately.
- dgrad with stride > 2, and direct dgrad with an odd number of
16-channel dY tiles (80, 112, …), are not yet cached.

JIRA-ID : AICK-2519

---------

Co-authored-by: Claude Opus 5.5 <noreply@anthropic.com>
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

Projects

None yet

Development

Successfully merging this pull request may close these issues.

3 participants