Context Log

2026-02-24 - 初回設計ドラフト作成

決定事項

まずは実装前に設計ドキュメントを整備する
cozipは「CPU + GPU(WebGPU)協調」を前提にチャンク並列設計を採用する
Deflateはチャンク独立性を高めるため、チャンク跨ぎ参照を禁止する方針を採用する
ZIP互換は維持しつつ、並列解凍用索引は extra field で保持する

採用理由

GPUへ割り当てるタスクを独立化しやすい
CPU/GPUの同時実行で総スループットを上げやすい
標準ZIPとの互換を維持しつつ、cozip同士で高速解凍経路を持てる

トレードオフ

圧縮率が通常Deflateより低下する可能性がある
小さな入力ではGPUオーバーヘッドが勝つ可能性がある

次アクション

Rustのモジュール設計(chunk, scheduler, deflate, zip, gpu)を src/ に反映
M1としてCPUのみのチャンク独立Deflate(最小実装)を作る
ZIPエンコード/デコードの基礎(ローカルヘッダ + セントラルディレクトリ)を追加

2026-02-24 - workspace分割 + Deflate実装着手

決定事項

ルート cozip を Cargo workspace 化し、cozip_deflate と cozip_zip に分割した
cozip_deflate に以下を実装した
純CPUの raw Deflate 圧縮・解凍関数
チャンク独立フォーマット(CZDF)による並列圧縮・解凍
CPU/GPU 協調実行(圧縮: GPU解析 + CPU Deflate、解凍: GPU解析支援 + CPU Deflate復元)
cozip_zip に最小の単一ファイル ZIP 圧縮・解凍を実装した

補足

現在のGPU経路は WebGPU でチャンク統計解析を実行し、圧縮レベル調整に利用している
Deflate のビットストリーム生成/復元自体は現段階ではCPU実装

次アクション

GPU側でLZ候補探索を持てるように cozip_deflate を段階拡張する
CZDFメタデータをZIP extra fieldに載せる統合を cozip_zip へ追加する
外部Deflate/ZIP実装との相互運用テストを増やす

2026-02-24 - CPU vs CPU+GPU 統合テスト追加

決定事項

cozip_deflate/tests/hybrid_integration.rs を追加
-- --nocapture 前提で CPU-only と CPU+GPU の比較ログを出力する
比較項目は圧縮/解凍時間、圧縮後サイズ、チャンク配分(stats)とした

仕様

GPU利用可能なら gpu_chunks > 0 を必須化
GPUがない環境ではフォールバック実行としてテスト継続し、ログで明示
追加テスト:

compare_cpu_only_vs_cpu_gpu_with_nocapture
hybrid_uses_both_cpu_and_gpu_when_gpu_is_available

実行コマンド

cargo test -p cozip_deflate --test hybrid_integration -- --nocapture
cargo test --workspace

2026-02-24 - GPU実タスク設計(大チャンク転送 + GPU内細分化)

決定事項

GPUは解析補助ではなく、Deflate圧縮/解凍の実タスクを担当する設計へ進める
2段階チャンク化を採用する
host_chunk(1〜4MiB, 初期2MiB): CPU/GPUスケジューリング単位
gpu_subchunk(64〜256KiB, 初期128KiB): GPU内部並列実行単位
静的50:50配分ではなく、共通キュー + EWMAによる動的配分を採用する

採用理由

転送オーバーヘッドを吸収しつつGPU並列度を確保できる
CPU/GPUの処理能力差があっても、先に空いた側へ仕事を回せる
可変長出力の競合を2パス(長さ計算->prefix sum->emit)で管理しやすい

追記したドキュメント

docs/gpu-full-task-design.md
docs/architecture.md (v1 draftへ更新)
docs/deflate-parallel-profile.md (host/subchunk方針へ更新)

次アクション

cozip_deflate に GpuContext 使い回し機構を導入する
HybridScheduler を共通キュー + 動的配分に置き換える
GPU圧縮本体(match_find, token_count, prefix_sum, token_emit)を段階実装する

2026-02-24 - CPU全力 + GPU全力実装(第1段)

決定事項

cozip_deflate を動的スケジューラへ置き換えた
CPUワーカー群とGPUワーカーが同時にキューを消費する実装にした
gpu_fraction に基づくGPU予約チャンクを導入し、CPUが取り切らないようにした
GpuContext はプロセス内で使い回す方式(遅延初期化)へ変更した

GPU実タスク(圧縮/解凍)

圧縮時:

GPUで連続一致率(repeat ratio)を計算し圧縮レベルを調整
GPUで EvenOdd 可逆変換(サブチャンク単位)を実行
CPUでDeflateビットストリーム化

解凍時:

CPUでDeflate展開
GPU(またはCPUフォールバック)で EvenOdd 逆変換

互換性

フレームバージョンを CZDF v2 へ更新
チャンクメタデータに transform フィールドを追加
旧 v1 フレームの読み取りは継続対応

検証

cargo test --workspace 通過
cargo test -p cozip_deflate --test hybrid_integration -- --nocapture で CPU と CPU+GPU 比較出力を確認
GPU利用可能環境で cpu_chunks / gpu_chunks 両方が配分されることを確認

2026-02-24 - 速度ベンチマーク実施(Release)

実行条件

コマンド: cargo run --release -p cozip_deflate --example bench_hybrid
反復回数: 5
データサイズ: 4MiB / 16MiB
方式比較:

CPU_ONLY (prefer_gpu=false, gpu_fraction=0.0)
CPU+GPU (prefer_gpu=true, gpu_fraction=0.5)

計測結果

4MiB:

CPU_ONLY: comp 72.977ms / decomp 4.276ms / comp 54.81MiB/s / decomp 935.43MiB/s
CPU+GPU : comp 156.421ms / decomp 6.821ms / comp 25.57MiB/s / decomp 586.41MiB/s

16MiB:

CPU_ONLY: comp 75.138ms / decomp 4.633ms / comp 212.94MiB/s / decomp 3453.26MiB/s
CPU+GPU : comp 682.850ms / decomp 7.577ms / comp 23.43MiB/s / decomp 2111.67MiB/s

所見

現状実装では CPU+GPU が CPU_ONLY を上回っていない
主因は GPU側タスクが「可逆変換 + 解析」に寄っており、Deflateビットストリーム本体がCPU実装のため
次段でGPU側のトークン生成本体(match_find/token_count/prefix_sum/token_emit)を強化する必要がある

2026-02-24 - 1GB比較用ベンチ追加

決定事項

cozip_deflate/examples/bench_1gb.rs を追加した
デフォルトを 1GiB 入力(--size-mib 1024)に設定した
CPU_ONLY と CPU+GPU を同条件で比較出力する
引数でサイズ/反復/ウォームアップ/チャンク設定を変更可能にした

実行コマンド

1GiB本番比較: cargo run --release -p cozip_deflate --example bench_1gb
軽量確認: cargo run --release -p cozip_deflate --example bench_1gb -- --size-mib 64 --iters 1 --warmups 0

次アクション

1GiB実測値を取得して共有
GPU側Deflate本体実装(match_find/token_count/prefix_sum/token_emit)へ着手

2026-02-24 - 1GBベンチのデフォルト実行回数を短縮

決定事項

bench_1gb のデフォルト反復数を短縮した
iters: 3 -> 1
warmups: 1 -> 0

理由

1GiB比較は1回でも十分傾向確認でき、待ち時間を大幅に減らせるため

補足

厳密比較したい場合は明示的に --iters / --warmups を指定する

2026-02-24 - GPU Deflate本体設計(独立チャンク実行 + 連結)

決定事項

実装方針を明確化した:

入力を独立チャンクへ分割
CPUとGPUがそれぞれDeflateを独立実行
index 順に戻して連結

圧縮率低下を許容し、並列スループットを優先する
Chunk-Member Profile (CMP) を採用し、各チャンクを独立Deflate memberとして扱う

追記したドキュメント

docs/gpu-deflate-chunk-pipeline.md
docs/architecture.md (CMPを追記)
docs/deflate-parallel-profile.md (CMPを追記)

次アクション

cozip_deflate に ChunkMember データモデルを固定する
GPU圧縮本体(match_find/token_count/prefix_sum/token_emit)を実装する
GPU解凍(固定Huffman先行)を追加する

2026-02-24 - GPU Deflate本体実装(第2段)

実装内容

cozip_deflate/src/lib.rs をCMP前提で更新した
GPU経路に match_find -> token_count -> prefix_sum -> token_emit を実装した
GPUで得た run 開始位置から、固定Huffman Deflate を生成する経路を追加した
各チャンクは独立 Deflate member として圧縮され、index 順で連結される

主要ポイント

GPU圧縮:

run_start_positions() でGPUパイプライン実行
encode_deflate_fixed_from_runs() で固定Huffman Deflate生成

CPU圧縮:

従来どおり flate2 Deflate

解凍:

チャンク単位でCPU/GPUワーカーが分担
Deflate展開は現段階ではCPU実装を利用
GPU担当チャンクはGPUパイプラインを補助的に実行可能

検証

2026-02-27 - D0実装（CZDIメタデータ基盤）

実装内容

cozip_deflate に DeflateChunkIndex / DeflateChunkIndexEntry を追加
CZDI v1 の encode/decode を実装（varint table + CRC）
圧縮APIを拡張し、deflate_compress_stream_zip_compatible_with_index で索引を返せるようにした
圧縮時のビット書き込み位置を追跡し、チャンクごとの comp_bit_off / comp_bit_len / final_header_rel_bit / raw_len を収集

ZIP統合

cozip で CZDIメタデータを書き込む実装を追加
- まず Central Directory extra field へ inline 格納
- 容量超過時は ZIP64 EOCD extensible data へ退避し、extraには locator を記録
読み取り側でも CZDI extra + EOCD64退避データを復元し、index decode まで実装

補足

D0段階のため、解凍実行はまだ既存CPU経路のまま（indexは読み取り・保持まで）
CoZipDeflate では未対応GPU解凍時にフォールバックせずエラー返却する方針を維持（実行切替は cozip 側責務）

cargo test -p cozip_deflate 通過
cargo test --workspace 通過
cargo test -p cozip_deflate --test hybrid_integration -- --nocapture 通過

補足

現在のGPU Deflate経路は「runベース + 固定Huffman」のため、圧縮率はCPU経路より悪化しやすい
次段はGPU側のトークン品質(一般LZ候補)と bitpack の拡張が必要

2026-02-24 - 改善1/3実装 (GPUトークン品質 + 大きめチャンク)

実装内容

改善1: GPUトークン品質向上

build_tokens_from_run_starts() を run-only 方式から LZ77 greedy 方式へ更新
ハッシュ候補 + runヒントの複数距離候補を評価し、最長一致を採用
既存GPUパイプライン(match_find/token_count/prefix_sum/token_emit)の出力をヒントとして利用

改善3: チャンクサイズのデフォルト拡大

HybridOptions::default():
chunk_size: 2MiB -> 4MiB
gpu_subchunk_size: 128KiB -> 256KiB
ベンチ既定値も同様に更新

計測(bench_hybrid, release, GPU実機)

4MiB:

CPU_ONLY ratio=0.3361
CPU+GPU ratio=0.3564

16MiB:

CPU_ONLY ratio=0.3364
CPU+GPU ratio=0.3465

所見

以前のGPU比率(約0.512)から大幅に改善
ただし速度は依然CPU_ONLY優位で、GPU bitpack/候補品質の追加改善が必要

追記テンプレート

## YYYY-MM-DD - タイトル
### 決定事項
- ...
### 問題
- ...
### 対応
- ...
### 次アクション
1. ...

2026-02-24 - 改善2/3着手 (GPUバッチsubmit + 転送前処理削減)

実装内容

改善3(実行バッチ化): GPU圧縮ワーカーを単一チャンク処理からバッチ処理へ変更

compress_gpu_worker() で最大 GPU_BATCH_CHUNKS 件(既定8)をまとめて取得
compress_chunk_gpu_batch() を追加し、複数チャンクのGPU補助処理を一括実行

改善2(転送・同期オーバーヘッド削減): run_start_positions_batch() を追加

旧 run_start_positions() を単発ラッパにし、内部はバッチAPI経由へ統一
複数チャンクの compute pass を1つの command encoder に積み、queue.submit() をバッチ単位に削減
map_async + poll(wait) もバッチ単位に集約し、同期回数を削減
GPU入力を u32 への1byte展開を廃止し、packed byte入力をWGSL側で参照する方式へ変更

検証

cargo test -p cozip_deflate 通過
cargo run --release -p cozip_deflate --example bench_hybrid 実行
cargo test --workspace 通過

メモ

小規模ベンチ(4MiB/16MiB)では改善幅が見えづらい
効果確認は1GiB級で再評価するのが妥当

2026-02-24 - GPU Deflate完結経路(固定Huffman literal) + readback最小化

実装内容

GPU内Deflate完結経路を追加

GpuAssist::deflate_fixed_literals_batch() を追加
GPU上で以下を実行して、チャンクごとのDeflateバイト列を生成
1. literal code/bitlen生成 (litlen_pipeline)
2. bit offset prefix-sum (prefix_pipeline 再利用)
3. bitpack (bitpack_pipeline)
CPU側は中間トークンを組み立てず、GPU出力をそのままチャンク圧縮データとして採用

readback最小化

旧経路の positions 大量readbackを圧縮経路から除去
圧縮で戻すデータを「合計bit数 + 最終圧縮バイト列」に限定

圧縮ワーカー接続変更

compress_chunk_gpu_batch() は run_start_positions_batch() + CPU bitpack ではなく、 deflate_fixed_literals_batch() を使用するよう変更

検証

cargo test -p cozip_deflate 通過
cargo test --workspace 通過
cargo run --release -p cozip_deflate --example bench_hybrid 実行

直近ベンチ(bench_hybrid)

4MiB:
- CPU_ONLY comp_mib_s=38.82 ratio=0.3361
- CPU+GPU comp_mib_s=22.95 ratio=1.0181
16MiB:
- CPU_ONLY comp_mib_s=122.75 ratio=0.3364
- CPU+GPU comp_mib_s=48.84 ratio=0.6771

メモ

速度面は以前のGPU経路より改善傾向だが、CPU_ONLYをまだ下回る
現在のGPU完結経路はliteral主体のため圧縮率が悪化しやすい
次段はGPU側でmatch探索/トークン化を強化し、match token(長さ・距離)を実際に出力する必要がある

2026-02-24 - 追加改善(継続): 安全側の高速化調整

対応

解凍時の不要GPU補助を削除

decode_descriptor_gpu() で行っていた run_start_positions() 呼び出しを削除
展開後データに対する補助GPU実行は、現行仕様では性能メリットが薄くオーバーヘッドが勝るため無効化

GPUバッチ粒度の拡大

GPU_BATCH_CHUNKS: 8 -> 16
1GiBクラスでGPU submit回数を減らしやすくする調整

重要メモ

共通バッファ再利用 + 単一 queue.submit() の試行は一旦見送り
理由: queue.write_buffer() は encoder内コマンドではないため、複数チャンク分の更新を先に積むと全dispatchが最終書き込み状態を参照し、データ破損が発生し得る
次に同方針を進める場合は、copy_buffer_to_buffer をencoder内に入れる staged upload 設計が必要

検証

cargo test -p cozip_deflate 通過
cargo test -p cozip_deflate --test hybrid_integration 通過
cargo run --release -p cozip_deflate --example bench_hybrid 実行

2026-02-24 - GPU len/dist 実装 (RLE系 match token 追加)

実装内容

GPUトークン化パスを追加

tokenize_pipeline を追加し、各バイトを literal / match(dist=1) / skip に分類
run先頭バイトは常に literal にし、残りを match(len, dist=1) 化

GPUコード生成を拡張

既存 litlen_pipeline をトークン入力ベースに変更
literal と match(len/dist) の固定HuffmanコードをGPU上で生成
bit長配列を作成してprefix-sum -> bitpack で最終Deflateを構築

readback最小化は維持

返却は total_bits + compressed bytes のみ

バグ修正

初版で dist=1 match開始位置が不正（run先頭でmatch開始）になり、展開破損
修正後は run先頭をliteral化し、残りのみmatch化して整合性を回復

検証

cargo test -p cozip_deflate 通過
cargo test -p cozip_deflate --test hybrid_integration -- --nocapture 通過
cargo test --workspace 通過
cargo run --release -p cozip_deflate --example bench_hybrid 実行

観測

圧縮率は literal-only より改善したケースがあるが、CPU_ONLYには依然届かない
現段階のmatchは dist=1 (RLE系) に限定されるため、一般データの圧縮効率はまだ不十分

2026-02-24 - GPU len/dist 強化 (candidate + finalize 2段)

実装内容

dist>1 を扱うGPU Deflateトークン化へ更新

token_dist バッファを追加し、match tokenに距離を保持
litlen シェーダで distance symbol/extra bits を生成するよう拡張

トークン生成を2段化

tokenize_pipeline:
- 各index独立で candidate(len, dist) を並列算出
- 候補距離は近距離優先の固定セット(1..1024の疎サンプル)
token_finalize_pipeline:
- 単一スレッドで候補列を走査し、非重複の最終token列へ確定
- literal / match(len,dist) / skip を整合性付きで確定

Deflate実行順を変更

tokenize(candidate) -> tokenize(finalize) -> token prefix -> litlen codegen -> bitlen prefix -> bitpack

目的と効果

以前の「各indexが独立にstart/skipを決める」方式で起きていた被覆競合を解消
圧縮長不整合(展開長ミスマッチ)を回避しつつ、GPUでlen/dist生成を継続

検証

cargo test -p cozip_deflate --lib 通過
cargo test -p cozip_deflate --test hybrid_integration -- --nocapture 通過

compare_cpu_only_vs_cpu_gpu_with_nocapture 成功
hybrid_uses_both_cpu_and_gpu_when_gpu_is_available 成功

現状メモ

correctnessは回復
速度はまだ CPU_ONLY より遅いが、GPUが空ブロック化する不安定挙動は解消
次段は tokenize_pipeline の候補探索をさらにGPUフレンドリー化(探索候補とscan長の最適化、shared memory活用)が必要

追加ベンチ (2026-02-24 / bench_hybrid)

4MiB
- CPU_ONLY: comp_mib_s=40.82 ratio=0.3361
- CPU+GPU : comp_mib_s=4.11 ratio=0.6592
16MiB
- CPU_ONLY: comp_mib_s=114.92 ratio=0.3364
- CPU+GPU : comp_mib_s=8.06 ratio=0.4977

メモ: correctnessは改善したが、性能・圧縮率ともCPU_ONLY優位のまま。

2026-02-24 - finalize/prefix ボトルネック削減

対応

token_finalize を単一点逐次からセグメント並列へ変更

旧: dispatch_workgroups(1,1,1) で全体を1スレッド処理
新: TOKEN_FINALIZE_SEGMENT_SIZE=4096 単位で複数workgroupへ分割
各セグメント内でgreedy確定(セグメント境界を跨がないよう match 長をクランプ)

prefix を階層並列scanへ置換

scan_blocks_pipeline で block 内 exclusive scan + block sums
block sums を再帰的に同scanで prefix 化
scan_add_pipeline で block offset を全要素へ加算
これを token prefix / bitlen prefix の両方で利用

実装メモ

追加定数: PREFIX_SCAN_BLOCK_SIZE=256, TOKEN_FINALIZE_SEGMENT_SIZE=4096
GpuAssist::dispatch_parallel_prefix_scan() を追加し、Deflate経路のprefix計算を統一

検証

cargo test -p cozip_deflate --lib 通過
cargo test -p cozip_deflate --test hybrid_integration -- --nocapture 通過
cargo run --release -p cozip_deflate --example bench_hybrid 実行

直近ベンチ (bench_hybrid)

4MiB
- CPU_ONLY: comp_mib_s=36.15 ratio=0.3361
- CPU+GPU : comp_mib_s=199.39 ratio=0.6593
16MiB
- CPU_ONLY: comp_mib_s=112.39 ratio=0.3364
- CPU+GPU : comp_mib_s=132.78 ratio=0.4978

補足: 圧縮スループットは大きく改善。圧縮率は依然CPU_ONLYより悪化しやすい。

2026-02-24 - Atomic状態ベースの適応スケジューラ

目的

固定比率(gpu_fraction)配分だけだとGPU遅い環境で全体が引きずられる問題を緩和
GPU優先予約を維持しつつ、未着手予約をCPUへ再配分する

実装

圧縮タスク状態をAtomic化

ScheduledCompressTask を追加
状態: Pending / ReservedGpu / RunningGpu / RunningCpu / Done
reserved_at_ms を保持し、GPU予約の鮮度を判定

GPU有効時の圧縮経路を新スケジューラへ切替

compress_hybrid() でGPU有効時は compress_hybrid_adaptive_scheduler() を使用
CPU-only時は既存キュー経路を維持

監視スレッド(Watchdog)を追加

ReservedGpu のまま一定時間(GPU_RESERVATION_TIMEOUT_MS)更新されないタスクを検出
CPU空き数(active_cpu)ぶんだけ Pending へ降格し、CPUに実行機会を渡す

CPU/GPUワーカーはCASでタスク獲得

CPUは Pending -> RunningCpu
GPUは ReservedGpu -> RunningGpu を優先、その後 Pending -> RunningGpu
実行後は Done に遷移し remaining を減算

待機は Condvar + timeout で実装

両者が取り合えない時は短時間sleepし、busy-spinを回避

検証

cargo test -p cozip_deflate --lib 通過
cargo test -p cozip_deflate --test hybrid_integration -- --nocapture 通過
cargo run --release -p cozip_deflate --example bench_hybrid 実行

直近ベンチ (bench_hybrid)

4MiB
- CPU_ONLY: comp_mib_s=40.85 ratio=0.3361
- CPU+GPU : comp_mib_s=190.71 ratio=0.6593
16MiB
- CPU_ONLY: comp_mib_s=138.54 ratio=0.3364
- CPU+GPU : comp_mib_s=149.78 ratio=0.4978

メモ: 圧縮率はまだGPU側が不利だが、スループット面は固定配分より改善。

追加ベンチ (2026-02-24 / 4GiB)

command: cargo run --release -p cozip_deflate --example bench_1gb -- --size-mib 4096 --iters 1 --warmups 0 --chunk-mib 4 --gpu-subchunk-kib 256

CPU_ONLY:
- comp_mib_s=514.72
- decomp_mib_s=4656.01
- ratio=0.3364
- cpu_chunks=1024 gpu_chunks=0
CPU+GPU:
- comp_mib_s=587.09
- decomp_mib_s=4282.01
- ratio=0.4020
- cpu_chunks=816 gpu_chunks=208

メモ:

圧縮はCPU+GPUが優位(+14%程度)
解凍はCPU_ONLYが優位
speedup表記の注記(>1.0 means CPU_ONLY faster)は逆で、実際はCPU_ONLY/hybrid比なので>1はhybridが速い

2026-02-24 - GPU転送/readback最適化 (実装4)

対応

不要なCPUゼロ書き込みを削減

deflate_fixed_literals_batch() で以下の queue.write_buffer(0埋め) を削除:
- token_total_buffer
- bitlens_buffer
- total_bits_buffer
- output_words_buffer の全体ゼロ埋め
output_words_buffer は先頭ヘッダ(0b011)のみ書き込み維持
WebGPUのゼロ初期化前提を利用し、PCIe転送量を削減

readbackを単一バッファ化

旧: total_readback と compressed_readback を別々にmap
新: readback 1本に total_bits(先頭4byte) + compressed payload を集約
map_async/map回数・チャネル処理を削減

検証

cargo test -p cozip_deflate --lib 通過
cargo test -p cozip_deflate --test hybrid_integration -- --nocapture 通過
cargo run --release -p cozip_deflate --example bench_1gb -- --size-mib 1024 --iters 1 --warmups 0 --chunk-mib 4 --gpu-subchunk-kib 256
cargo run --release -p cozip_deflate --example bench_hybrid

直近ベンチ

1GiB (bench_1gb)
- CPU_ONLY: comp_mib_s=476.38 decomp_mib_s=4146.80 ratio=0.3364
- CPU+GPU : comp_mib_s=572.34 decomp_mib_s=5232.33 ratio=0.4373
- chunk配分: cpu=176 gpu=80
bench_hybrid
- 4MiB:
  - CPU_ONLY comp_mib_s=39.12 ratio=0.3361
  - CPU+GPU comp_mib_s=237.52 ratio=0.6593
- 16MiB:
  - CPU_ONLY comp_mib_s=95.98 ratio=0.3364
  - CPU+GPU comp_mib_s=155.51 ratio=0.4978

メモ

圧縮スループットは改善傾向
圧縮率は依然としてCPU_ONLYより悪化しやすく、GPU match品質の改善が次段課題

2026-02-24 - 速度最適化(1+2継続): Deflateスロット再利用

実装

GPU Deflateバッファの永続化/再利用

GpuAssist に deflate_slots: Mutex<Vec<DeflateSlot>> を追加
チャンクごとの大量 create_buffer/create_bind_group を削減
必要容量を超えた場合のみスロットを再確保

per-slot bind group 再利用

litlen_bg / tokenize_bg / bitpack_bg をスロット内に保持
毎チャンク再生成を廃止

初期化のGPU側クリア

再利用バッファの初期化は encoder.clear_buffer を利用
output_words ヘッダは専用 deflate_header_buffer から copy_buffer_to_buffer で設定

readbackは単一バッファ維持

total_bits + payload を1本で回収

検証

cargo test -p cozip_deflate --lib 通過
cargo test -p cozip_deflate --test hybrid_integration -- --nocapture 通過
cargo run --release -p cozip_deflate --example bench_1gb -- --size-mib 1024 --iters 1 --warmups 0 --chunk-mib 4 --gpu-subchunk-kib 256
cargo run --release -p cozip_deflate --example bench_hybrid

直近ベンチ

1GiB (bench_1gb)
- CPU_ONLY: comp_mib_s=502.32 decomp_mib_s=4364.46 ratio=0.3364
- CPU+GPU : comp_mib_s=625.10 decomp_mib_s=5463.49 ratio=0.4373
- speedup(cpu/hybrid): compress=1.244x decompress=1.252x
- 配分: cpu_chunks=176 gpu_chunks=80
bench_hybrid
- 4MiB:
  - CPU_ONLY comp_mib_s=33.06 ratio=0.3361
  - CPU+GPU comp_mib_s=305.02 ratio=0.6593
- 16MiB:
  - CPU_ONLY comp_mib_s=112.70 ratio=0.3364
  - CPU+GPU comp_mib_s=133.84 ratio=0.4978

メモ

1 は反映済み
2 の完全版(二重バッファ submit/collect 分離による upload/compute/readback 重畳)は次段実装候補

2026-02-24 - 速度最適化(2完全版): submit/collect 重畳の実装

実装

deflate_fixed_literals_batch() をスロットプール型に変更

chunk_index 固定スロットではなく、free_slots + pending で再利用する方式へ移行
プール上限: GPU_DEFLATE_SLOT_POOL (= GPU_BATCH_CHUNKS)

submit と collect の重畳

GPU_PIPELINED_SUBMIT_CHUNKS ごとに queue.submit して map_async 登録
直後に poll(Maintain::Poll) + try_recv で回収可能分を先行 collect
空きスロットがない場合のみ poll(Maintain::Wait) で1件回収して再利用

readback復元処理の共通化

collect_deflate_readback() を追加
total_bits + payload 形式の readback から圧縮バイト列を復元し、unmap まで一元化

検証

cargo test -p cozip_deflate --lib 通過
cargo test -p cozip_deflate --test hybrid_integration -- --nocapture 通過
cargo run --release -p cozip_deflate --example bench_1gb 実行

直近ベンチ

1GiB (bench_1gb, default)
- CPU_ONLY: comp_mib_s=499.68 decomp_mib_s=4031.88 ratio=0.3364
- CPU+GPU : comp_mib_s=635.92 decomp_mib_s=5209.82 ratio=0.4574
- speedup(cpu/hybrid): compress=1.273x decompress=1.292x
- 配分: cpu_chunks=160 gpu_chunks=96

メモ

前回実装の「部分的重畳」から、実際に submit/collect を分離したパイプラインへ移行
圧縮率は CPU_ONLY より悪いままだが、スループットは引き続き CPU+GPU が優位

2026-02-24 - モード共存実装 (Speed/Balanced/Ratio + codec_id)

実装

圧縮モードを追加

HybridOptions に compression_mode: CompressionMode を追加
CompressionMode:
- Speed (既存優先)
- Balanced (GPU探索品質を中間設定)
- Ratio (GPU探索品質を高設定)
Default は CompressionMode::Speed

フレームメタへ codec_id を追加

FRAME_VERSION を 3 へ更新
chunk metadata に codec_id を追加 (backend + transform + codec + raw_len + compressed_len)
新形式書き出し時は v3
読み取りは v1/v2/v3 互換維持
- v2は backend から codec を推定

チャンクcodec分岐を追加

ChunkCodec::{DeflateCpu, DeflateGpuFast} を導入
圧縮時にチャンクごとに codec を格納
復号時は codec で分岐してinflate（現状どちらも deflate inflate だが将来拡張点を確保）

Balanced / Ratio の挙動

speed と同様に GPU へタスクを割り当てる
CPU再圧縮フォールバックではなく、GPU tokenize/finalize の探索品質をモード別に切り替える

検証

cargo test -p cozip_deflate --lib 通過
cargo test -p cozip_deflate --test hybrid_integration -- --nocapture 通過
cargo test -p cozip_deflate 通過
追加テスト:

ratio_mode_roundtrip
decode_v2_frame_compatibility

運用メモ

bench_1gb に --mode speed|balanced|ratio を追加済み

4GiB モード比較 (2026-02-24 / ローカル実測・更新版)

command (各モード): target/release/examples/bench_1gb --size-mib 4096 --iters 1 --warmups 0 --chunk-mib 4 --gpu-subchunk-kib 256 --mode <speed|balanced|ratio>

結果:

speed
- CPU_ONLY: comp_mib_s=514.68 decomp_mib_s=4328.25 ratio=0.3364
- CPU+GPU : comp_mib_s=753.17 decomp_mib_s=4888.60 ratio=0.4625
- speedup(cpu/hybrid): compress=1.463x decompress=1.129x
- chunk配分: cpu_chunks=624 gpu_chunks=400
- GPU使用率観測: ほぼ100%張り付き
balanced
- CPU_ONLY: comp_mib_s=518.28 decomp_mib_s=4339.05 ratio=0.3364
- CPU+GPU : comp_mib_s=435.05 decomp_mib_s=5488.71 ratio=0.3364
- speedup(cpu/hybrid): compress=0.839x decompress=1.265x
- chunk配分: cpu_chunks=1024 gpu_chunks=0
- GPU使用率観測: 16%前後
ratio
- CPU_ONLY: comp_mib_s=514.65 decomp_mib_s=4328.79 ratio=0.3364
- CPU+GPU : comp_mib_s=527.86 decomp_mib_s=3767.50 ratio=0.3364
- speedup(cpu/hybrid): compress=1.026x decompress=0.870x
- chunk配分: cpu_chunks=1024 gpu_chunks=0
- GPU使用率観測: 30%前後

所見:

圧縮速度最優先なら speed が最良。CPU+GPUで圧縮/解凍ともCPU_ONLYを上回る。
圧縮率最優先なら balanced が有効（CPU_ONLY同等の ratio を維持）。ただし圧縮スループットは低下。
ratio は現状ほぼCPU実行だが、圧縮速度はCPU_ONLYと同等以上、解凍速度は低下傾向が見られる。

2026-02-24 - モード仕様修正 (GPU再計算フォールバック廃止)

方針:

balanced / ratio でも speed と同様に GPU へ圧縮タスクを割当
圧縮率改善は CPU再圧縮でなく GPU側ロジック改善で行う

実装:

compress_hybrid() の Ratio による GPU無効化を廃止
GPU圧縮の CPU再圧縮比較ロジックを撤去
tokenize shader:
- mode別に max_match_scan / max_match_len / distance candidate count を切替
- distance candidate を 32段階 (最大 32768) まで拡張
token_finalize shader:
- mode別 lazy matching (speed=0, balanced=1, ratio=2) を追加

mode別GPU品質パラメータ:

Speed: scan=64, max_len=64, dist_slots=20
Balanced: scan=128, max_len=128, dist_slots=28
Ratio: scan=192, max_len=258, dist_slots=32

検証:

cargo test -p cozip_deflate --lib 通過
cargo test -p cozip_deflate --test hybrid_integration -- --nocapture 通過

2026-02-24 - Ratioモード: GPU頻度集計 + CPU木生成(dynamic Huffman)

実装:

ratio モードでも GPU タスク割当は維持（GPU無効化しない）
deflate_fixed_literals_batch() に ratio 分岐を追加し、deflate_dynamic_hybrid_batch() を使用
deflate_dynamic_hybrid_batch():
1. GPUで tokenize + finalize
2. GPU frequency pass で litlen(286) / dist(30) 頻度を atomic 集計
3. 頻度テーブルを readback
4. CPUで Huffman木(符号長)生成 + canonical code生成
5. CPUで dynamic header + token列を bitpack
GPU側は mode に応じて探索品質を切替:
- Speed: scan=64 / len=64 / dist_slots=20
- Balanced: scan=128 / len=128 / dist_slots=28
- Ratio: scan=192 / len=258 / dist_slots=32

補足:

balanced/ratio のCPU再圧縮フォールバックは廃止
圧縮率改善はスケジューリングではなく GPU tokenization 品質 + dynamic Huffman で行う

検証:

cargo test -p cozip_deflate --lib 通過
cargo test -p cozip_deflate --test hybrid_integration -- --nocapture 通過
cargo check -p cozip_deflate --examples 通過

2026-02-24 - Balanced/Ratio 実行時エラー修正

報告エラー:

mode=balanced で InvalidFrame("raw chunk length mismatch in cpu path")
mode=ratio で Source buffer is missing the COPY_SRC usage flag

修正:

ratio用 dynamic path で readback コピーしているバッファに COPY_SRC を追加
- token_flags/token_kind/token_len/token_dist/token_lit
balanced/ratio の GPU 圧縮結果に対して整合性ガードを追加
- GPU圧縮結果を inflate して元チャンク一致を検証
- 不一致時のみ CPU圧縮へフォールバック（クラッシュ防止の保険）

検証:

cargo test -p cozip_deflate --lib 通過
cargo test -p cozip_deflate --test hybrid_integration -- --nocapture 通過

2026-02-24 - Ratio 1+2 実装 (freq+final readback / GPU bitpack)

目的:

ratio で readback を最小化
- 旧: token配列 + freq + 最終出力
- 新: freq + 最終出力のみ
CPUは dynamic Huffman 木生成のみ
GPUで token bitpack + EOB finalize を実行

主な変更:

deflate_dynamic_hybrid_batch() を2段化
1. GPU tokenize/finalize/freq → litlen/dist 頻度のみreadback
2. CPUで dynamic Huffman plan 作成
3. planをGPUへupload
4. GPUで token map → prefix scan → bitpack → dynamic finalize(EOB)
5. 最終圧縮バイト列のみreadback
dynamic Huffman plan構築ヘルパーを追加
- build_dynamic_huffman_plan()
GPU dynamic map/finalize パイプラインを追加
- dyn_map_pipeline
- dyn_finalize_pipeline
bitpack shader を拡張
- base bit offset を params._pad1 で可変化
- 33bit以上コード用の high-lane (dyn_overflow_buffer) を追加

制約対応:

wgpu の max_storage_buffers_per_shader_stage=8 制限へ対応
- dyn_map のstorage bindingを8本以内に再設計
- token compact index(token_prefix)依存を除去（lane indexで直接bitpack）
- dynamic tableを単一storage buffer (dyn_table_buffer) に統合

バッファ/slot側:

追加: dyn_table_buffer, dyn_meta_buffer, dyn_overflow_buffer, dyn_map_bg, dyn_finalize_bg
出力上限を dynamic も見込んで拡張
- GPU_DEFLATE_MAX_BITS_PER_BYTE = 20

検証:

cargo test -p cozip_deflate --lib --no-run 通過
cargo test -p cozip_deflate --test hybrid_integration -- --nocapture 通過
cargo run --release -p cozip_deflate --example bench_1gb -- --size-mib 512 --iters 1 --warmups 0 --chunk-mib 4 --gpu-subchunk-kib 256 --mode ratio 通過
- CPU_ONLY ratio=0.3364, CPU+GPU ratio=0.3373 (512MiB, 1iter)

備考:

未使用関数/定数に関する warning は残る（既存設計由来）。
今回は panic/validation error を出さずに ratio 経路をGPU bitpack 化できた状態。

2026-02-24 - bench_1gb: gpu_fraction引数追加 + ratio 4GiB試験(1.0)

変更:

cozip_deflate/examples/bench_1gb.rs に --gpu-fraction <F> を追加
CPU+GPUケースの HybridOptions.gpu_fraction をCLI指定値で上書き
ベンチ出力に gpu_fraction を表示

実行:

cargo run --release -p cozip_deflate --example bench_1gb -- --size-mib 4096 --iters 1 --warmups 0 --chunk-mib 4 --gpu-subchunk-kib 256 --mode ratio --gpu-fraction 1.0

結果:

CPU_ONLY: comp=518.47 MiB/s, decomp=4412.30 MiB/s, ratio=0.3364
CPU+GPU : comp=583.18 MiB/s, decomp=5803.31 MiB/s, ratio=0.3375
speedup(cpu/hybrid): compress=1.125x, decompress=1.315x
chunk配分: cpu_chunks=833, gpu_chunks=191

所見:

gpu_fraction=1.0 でも実配分は動的調整でCPU側へ多く戻る（予約比率=実配分ではない）。
圧縮率差は小さいまま、圧縮/解凍ともCPU+GPUが優位。

2026-02-24 - デフォルトgpu_fractionを1.0へ変更

変更:

HybridOptions::default().gpu_fraction を 0.5 -> 1.0 に変更
bench_1gb の --gpu-fraction デフォルト表示/実値を 1.0 に変更

確認:

cargo test -p cozip_deflate --lib 通過

2026-02-24 - balanced低GPU利用の原因調査と修正

現象:

--mode balanced で gpu_chunks が極端に少ない / 0 になるケースが発生
圧縮速度がCPU_ONLYより悪化

原因(確認済み):

dynamic Huffman計画の code-length 木生成失敗

エラー: failed to build codelen huffman lengths
これにより GPUバッチ全体がCPUフォールバックし、gpu_chunks が増えない

予約降格タイミングが短く、GPU予約が早期にCPUへ流れる

旧設計は予約時刻が同時刻で、watchdogがまとめて降格しやすかった

実装修正:

dynamic Huffman code-length木生成にフォールバックを追加
- fallback_codelen_lengths()
- 生成不能時は code-lengthシンボルに安全な固定長(5bit)を割当
予約降格のモード別チューニング
- GPU_RESERVATION_TIMEOUT_MS_DYNAMIC = 100
- GPU_RESERVATION_STAGGER_MS_DYNAMIC = 8
予約時刻を段階的にずらす初期化を追加
- reserved_at_ms = now + seq * stagger_ms
- 一斉降格を抑止し、GPUが連続してバッチ取得しやすくした
balanced は引き続き dynamic Huffman + speed探索（tokenize modeはSpeed）
GPU検証は balanced/ratio で有効維持（誤圧縮防止）

補助デバッグ:

COZIP_LOG_GPU_FALLBACK=1 で以下をstderr出力
- GPUバッチエラー
- GPUバッチサイズ不整合
- GPU検証失敗によるCPUフォールバック

ローカル確認(1GiB, balanced, gpu_fraction=1.0):

修正前: gpu_chunks=0 相当のケースあり
修正後例: gpu_chunks=72 / cpu_chunks=184
ただし圧縮率はまだ高め (ratio=0.3967 例) で、balancedの圧縮品質は引き続き改善余地あり

2026-02-24 - gpu-fractionフラグ再追加（再適用）

対応内容:

cozip_deflate/examples/bench_1gb.rs
- --gpu-fraction <R> を再追加（0.0..=1.0）
- デフォルトを 1.0 に設定
- CPU+GPU 実行時の HybridOptions.gpu_fraction に反映
- ベンチ出力に gpu_fraction=... を表示
cozip_deflate/src/lib.rs
- HybridOptions::default().gpu_fraction を 1.0 に変更

確認:

cargo check -p cozip_deflate --example bench_1gb 通過

2026-02-24 - GPU検出確認用 nocapture テスト追加

目的:

cargo test ... -- --nocapture で、現在環境で見えているGPUと、PowerPreference別に選択されるGPUを確認できるようにする。

変更:

cozip_deflate/tests/hybrid_integration.rs
- print_current_gpu_with_nocapture テストを追加
- 検出アダプタ一覧（name/vendor/device/type/backend/driver）を表示
- HighPerformance / LowPower それぞれで request_adapter 結果を表示
- それぞれ request_device の成否を表示
- GPU未検出でも panic せずに情報出力して通る

ローカル実行結果:

検出:
- NVIDIA GeForce RTX 5070 Laptop GPU (Vulkan)
- AMD Radeon 890M Graphics (GL)
選択:
- HighPerformance: NVIDIA
- LowPower: NVIDIA

2026-02-24 - GPU dispatch 2D化 + submit/collect待機点削減 (着手: 1,2)

目的:

16MiB級チャンクで dispatch_workgroups(x>65535) に当たる問題を解消
GPU圧縮バッチで待機点を減らし、submit/collectをより分離

実装:

2D dispatch 対応

追加:
- MAX_DISPATCH_WORKGROUPS_PER_DIM = 65535
- dispatch_grid_for_groups()
- dispatch_grid_for_items()
変更:
- pass.dispatch_workgroups(..) を主要GPUパスで (x,y,1) に変更
- 対象: tokenize/litlen/bitpack/freq/dyn_map/scan blocks/scan add/match/count/emit
WGSL側 index を 2D flatten 対応
- workgroup_size=128 系: idx = id.x + id.y * 8388480
- workgroup_size=256 系: idx = id.x + id.y * 16776960
- scan_blocks は gid = wg_id.x + wg_id.y * 65535
scan_blocks の over-dispatch 対策
- block_sums[gid] 書き込み時に gid*256 < len ガード追加

submit/collect の待機点削減（fixed GPU batch path）

deflate_fixed_literals_batch で、slot枯渇時に Wait 後まとめて ready readback を回収
最終回収で poll(Wait) を1回化し、pendingを連続 drain
1件ずつ Wait していた箇所を削減

確認:

cargo check -p cozip_deflate 通過
cargo test -p cozip_deflate --lib 通過
cargo test -p cozip_deflate --test hybrid_integration -- --nocapture 通過
cargo run --release -p cozip_deflate --example bench_hybrid 実行
- 4MiB: CPU+GPU comp ~11.679ms
- 16MiB: CPU+GPU comp ~94.245ms
cargo run --release -p cozip_deflate --example bench_1gb -- --size-mib 16 --iters 1 --warmups 1 --chunk-mib 16 --gpu-subchunk-kib 256 --mode speed --gpu-fraction 1.0
- 16MiB単一チャンク (chunk_mib=16) が panicせず通ることを確認

メモ:

2D flatten の stride は global_invocation_id.x の性質上、workgroup_size込みで設定する必要がある。

2026-02-24 - dynamic側バッチ最適化（phase分離）

目的:

dynamic Huffman 経路のチャンク毎 submit->map->wait を削減し、GPU待機点を減らす。

変更:

deflate_dynamic_hybrid_batch を2段階バッチに再構成
1. Phase1: tokenize/finalize/freq を複数チャンクまとめて submit
  - 周期: GPU_PIPELINED_SUBMIT_CHUNKS
  - submit後に freq readback を map 予約し、最後に一括 poll(Wait) で回収
2. Phase2: CPUで Huffman plan 生成後、dyn_map/bitpack/finalize を複数チャンクまとめて submit
  - 同様に readback map を束ね、最後に一括回収
追加した内部構造体
- PendingDynFreqReadback
- PreparedDynamicPack
- PendingDynPackReadback

期待効果:

dynamic経路での同期オーバーヘッド低減
チャンク数が増えたときの submit/wait の効率化

確認:

cargo check -p cozip_deflate 通過
cargo test -p cozip_deflate --lib 通過
cargo test -p cozip_deflate --test hybrid_integration -- --nocapture 通過
cargo run --release -p cozip_deflate --example bench_1gb -- --size-mib 64 --iters 1 --warmups 1 --chunk-mib 4 --gpu-subchunk-kib 256 --mode ratio --gpu-fraction 1.0
- 実行成功（roundtrip OK）

2026-02-25 - GPU圧縮の時間計測ログ追加（COZIP_PROFILE_TIMING）

目的:

大幅改善余地の有無を切り分けるため、GPU圧縮パスのどこで時間を使っているかを可視化。

実装:

cozip_deflate/src/lib.rs
- COZIP_PROFILE_TIMING=1 で有効化される軽量タイミング計測を追加
- 追加ヘルパー:
  - timing_profile_enabled()
  - elapsed_ms()
  - GPU_TIMING_CALL_SEQ（ログ追跡用call id）
- fixed GPU path (deflate_fixed_literals_batch) サマリ出力:
  - t_encode_submit_ms
  - t_bits_rb_ms（total_bits readback待ち）
  - t_payload_submit_ms
  - t_payload_rb_ms
  - t_cpu_fallback_ms
  - payload_chunks/fallback_chunks/readback量
- dynamic GPU path (deflate_dynamic_hybrid_batch) サマリ出力:
  - t_freq_submit_ms
  - t_freq_wait_plan_ms（freq回収+CPU木生成）
  - t_pack_submit_ms
  - t_pack_bits_rb_ms（total_bits回収）
  - t_payload_submit_ms
  - t_payload_rb_ms
  - t_cpu_fallback_ms
  - payload_chunks/fallback_chunks/readback量
- adaptive scheduler (compress_hybrid_adaptive_scheduler) サマリ出力:
  - 総時間、CPU/GPUチャンク数、入出力サイズ

使い方:

例:
- COZIP_PROFILE_TIMING=1 cargo run --release -p cozip_deflate --example bench_1gb -- --size-mib 4096 --iters 1 --warmups 0 --chunk-mib 4 --gpu-subchunk-kib 256 --mode speed --gpu-fraction 1.0

ローカル動作確認:

cargo check -p cozip_deflate 通過
cargo check -p cozip_deflate --example bench_1gb 通過
COZIP_PROFILE_TIMING=1 cargo run --release -p cozip_deflate --example bench_1gb -- --size-mib 8 --iters 1 --warmups 0 --chunk-mib 4 --gpu-subchunk-kib 256 --mode speed --gpu-fraction 1.0
- [cozip][timing][gpu-fixed] ...
- [cozip][timing][scheduler] ... が出力されることを確認

2026-02-25 - dynamic freq区間の犯人切り分け（poll/map/plan分離）

目的:

t_freq_wait_plan_ms が大きい問題を、GPU待機かCPU木生成かまで分解して特定する。

実装:

GpuDynamicTiming を分割:
- t_freq_poll_wait_ms
- t_freq_recv_ms
- t_freq_map_copy_ms
- t_freq_plan_ms
deflate_dynamic_hybrid_batch で
- device.poll(Wait) 区間
- receiver recv 区間
- map + copy + unmap 区間
- build_dynamic_huffman_plan 区間を個別計測。

ローカル確認 (64MiB, ratio):

COZIP_PROFILE_TIMING=1 cargo run --release -p cozip_deflate --example bench_1gb -- --size-mib 64 --iters 1 --warmups 0 --chunk-mib 4 --gpu-subchunk-kib 512 --mode ratio --gpu-fraction 1.0
出力例:
- t_freq_poll_wait_ms=283.916
- t_freq_recv_ms=0.001
- t_freq_map_copy_ms=0.021
- t_freq_plan_ms=0.386

所見:

少なくともこの実行では、freq区間の大半は「GPU freqカーネル完了待ち (poll wait)」。
CPU側の木生成は支配的でない。

2026-02-25 - ratio: freq集計をworkgroup局所化 + capped dispatch

目的:

freq フェーズの global atomic 競合を下げ、t_freq_poll_wait_ms を短縮する。

実装:

cozip_deflate/src/lib.rs
- 追加: GPU_FREQ_MAX_WORKGROUPS = 4096
- 追加: dispatch_grid_for_items_capped(items, group_size, max_groups)
- dynamic の freq pass dispatch を
  - dispatch_grid_for_items(len, 128) から
  - dispatch_grid_for_items_capped(len, 128, GPU_FREQ_MAX_WORKGROUPS) に変更
- freq WGSL を変更:
  - 直接 litlen_freq/dist_freq へ atomicAdd する方式を廃止
  - var<workgroup> local_litlen_freq/local_dist_freq に集計
  - workgroup内で集計後、非0 binのみ global freq へ atomicAdd
  - grid-stride loop (idx += num_workgroups*workgroup_size) で1スレッドが複数トークン処理

確認:

cargo check -p cozip_deflate 通過
COZIP_PROFILE_TIMING=1 cargo run --release -p cozip_deflate --example bench_1gb -- --size-mib 64 --iters 1 --warmups 0 --chunk-mib 4 --gpu-subchunk-kib 512 --mode ratio --gpu-fraction 1.0
- ローカル結果:
  - t_freq_poll_wait_ms=269.699（前回計測 283.916 から減少）
  - t_freq_plan_ms=0.756（CPU木生成は依然支配的でない）

メモ:

改善幅は限定的で、さらに詰めるには GPU_FREQ_MAX_WORKGROUPS のチューニング、または partial histogram バッファを使った完全2pass reduce（global atomic最小化）が候補。

2026-02-25 - dynamic Phase1 深掘りプローブ追加（tokenize/finalize/freq分離）

目的:

Phase1(tokenize + token_finalize + freq)の真犯人を予測ではなく計測で特定する。

実装:

COZIP_PROFILE_DEEP=1 を追加（COZIP_PROFILE_TIMING と併用推奨）
GpuAssist::profile_dynamic_phase1_probe() を追加し、dynamic pathで最初の非空chunkに対して
- tokenize pass を単独submit+wait
- token_finalize pass を単独submit+wait
- freq pass を単独submit+wait を実行し、各msを出力
出力形式:
- [cozip][timing][gpu-dynamic-probe] ... t_tokenize_ms=... t_finalize_ms=... t_freq_ms=...

確認:

cargo check -p cozip_deflate 通過
COZIP_PROFILE_TIMING=1 COZIP_PROFILE_DEEP=1 cargo run --release -p cozip_deflate --example bench_1gb -- --size-mib 64 --iters 1 --warmups 0 --chunk-mib 4 --gpu-subchunk-kib 512 --mode ratio --gpu-fraction 1.0
ローカル例:
- t_tokenize_ms=17.369
- t_finalize_ms=2.386
- t_freq_ms=0.171

所見:

少なくともこの計測では、freq ではなく tokenize がPhase1支配要因。

2026-02-25 - tokenize内訳プローブ（literal/head/extend 差分計測）

目的:

tokenize 内の真犯人を特定するため、処理内訳を差分で計測。

実装:

tokenize WGSL に deep profile用 mode を追加
- 100: literal-only（候補探索なし）
- 101: head-only speed
- 102: head-only balanced
- 103: head-only ratio
profile_dynamic_phase1_probe() を拡張
- tokenize を 3回実行（lit/head/full）
- head_only = head_total - lit
- extend_only = full - head_total
- 各モードは warmup 1回 + 計測1回でブレを低減
出力形式:
- [cozip][timing][gpu-dynamic-probe] ... t_tokenize_lit_ms=... t_tokenize_head_total_ms=... t_tokenize_full_ms=... t_tokenize_head_only_ms=... t_tokenize_extend_only_ms=...

確認:

cargo check -p cozip_deflate 通過
COZIP_PROFILE_TIMING=1 COZIP_PROFILE_DEEP=1 cargo run --release -p cozip_deflate --example bench_1gb -- --size-mib 64 --iters 1 --warmups 0 --chunk-mib 4 --gpu-subchunk-kib 512 --mode ratio --gpu-fraction 1.0
ローカル例(4MiB probe):
- t_tokenize_lit_ms=0.313
- t_tokenize_head_only_ms=0.486
- t_tokenize_extend_only_ms=15.653

所見:

tokenize内では extend_only（一致後の長さ延長ループ）が圧倒的に支配的。

2026-02-25 - tokenize延長ループ最適化（4byte比較）

目的:

真犯人だった tokenize_extend_only を直接短縮する。

実装:

tokenize WGSL に load_u32_unaligned() を追加
一致後延長ループを
- 旧: 1byteずつ byte_at(p) == byte_at(p-dist) 比較
- 新: 4byte比較 (left4 == right4) を先行し、ミスマッチ時は countTrailingZeros(xor)>>3 で差分byte位置まで一気に進める
- 残りは tail の1byteループで処理
先頭3byte判定の byte_at(i), byte_at(i+1), byte_at(i+2) を事前読み出しして再利用

確認:

cargo check -p cozip_deflate 通過
COZIP_PROFILE_TIMING=1 cargo run --release -p cozip_deflate --example bench_1gb -- --size-mib 64 --iters 1 --warmups 0 --chunk-mib 4 --gpu-subchunk-kib 512 --mode ratio --gpu-fraction 1.0
- t_freq_poll_wait_ms: 283ms級 -> 149ms級（ローカル）
COZIP_PROFILE_TIMING=1 COZIP_PROFILE_DEEP=1 ... で深掘り
- t_tokenize_extend_only_ms: 15.6ms級 -> 7.3ms級（4MiB probe, ローカル）

メモ:

深掘り (COZIP_PROFILE_DEEP=1) は計測用追加実行が入るため、実運用ベンチ比較には使わない。

2026-02-25 - tokenize延長ループの追加最適化（16byte先行比較）

目的:

4byte比較化の次段として、長い一致ランでの反復回数をさらに削減する。

実装:

tokenize WGSL の延長ループに 16byte 先行比較を追加
- 4byte×4本の load_u32_unaligned 比較
- 途中不一致時は xor + countTrailingZeros >> 3 で一致バイト数を即時反映
- 16byte一致時のみ mlen/p/scanned を +16
- その後は既存の4byteループ -> 1byte tail へフォールバック

確認:

cargo check -p cozip_deflate 通過
COZIP_PROFILE_TIMING=1 COZIP_PROFILE_DEEP=1 cargo run --release -p cozip_deflate --example bench_1gb -- --size-mib 64 --iters 1 --warmups 0 --chunk-mib 4 --gpu-subchunk-kib 512 --mode ratio --gpu-fraction 1.0
- ローカル例(4MiB probe):
  - t_tokenize_extend_only_ms: 15.6ms級 -> 7.6ms級
  - t_tokenize_full_ms: 16.8ms級 -> 8.4ms級
- dynamic全体:
  - t_freq_poll_wait_ms: 280-313ms級 -> 148-170ms級

メモ:

COZIP_PROFILE_DEEP=1 実行では追加プローブ分のオーバーヘッドがあるため、最終スループット比較は COZIP_PROFILE_TIMING=1 のみで行う。

2026-02-25 - deep計測の誤比較防止

実装:

COZIP_PROFILE_DEEP 有効時に警告を1回だけ表示
- throughput numbers include probe overhead ...
deep probe 実行回数を「各dynamicバッチ」から「プロセス全体で1回」に変更
- DEEP_DYNAMIC_PROBE_TAKEN で制御
COZIP_PROFILE_TIMING=1 時に選択GPUアダプタ情報を出力
- name/vendor/device/backend/type

目的:

deep計測が有効なままベンチ比較してしまう事故を減らす
ハイブリッドGPU環境でアダプタ揺れを観測しやすくする

2026-02-25 - ベースライン再計測（ウォームアップなし / プロセス再起動）

目的:

zip実運用想定（事前ウォームアップなし）に合わせ、iters=1 warmups=0 で毎回プロセス起動し直して現状ベースラインを取得。

条件:

コマンド:
- env -u COZIP_PROFILE_DEEP -u COZIP_PROFILE_TIMING target/release/examples/bench_1gb --size-mib 4096 --iters 1 --warmups 0 --chunk-mib 4 --gpu-subchunk-kib 512 --mode ratio --gpu-fraction 1.0
実行回数: 5回（毎回プロセス再起動）

観測値:

Run1: CPU_ONLY comp=7898.563ms / CPU+GPU comp=19398.078ms / speedup(comp)=0.407x / speedup(decomp)=0.984x
Run2: CPU_ONLY comp=7893.687ms / CPU+GPU comp=19294.458ms / speedup(comp)=0.409x / speedup(decomp)=1.002x
Run3: CPU_ONLY comp=7996.905ms / CPU+GPU comp=19235.450ms / speedup(comp)=0.416x / speedup(decomp)=1.011x
Run4: CPU_ONLY comp=8005.602ms / CPU+GPU comp=19223.456ms / speedup(comp)=0.416x / speedup(decomp)=1.002x
Run5: CPU_ONLY comp=8011.265ms / CPU+GPU comp=19300.104ms / speedup(comp)=0.415x / speedup(decomp)=1.013x

集計:

speedup(comp) median=0.415x, mean=0.413x
speedup(decomp) median=1.002x, mean=1.002x

補足:

本実行環境では libEGL warning: failed to open /dev/dri/... Permission denied が発生。
gpu_chunks=16/1024 と極端に少なく、GPU性能比較としては無効寄り（GPUデバイス権限制約の影響）。
この記録は「手順の再現確認」と「無ウォームアップ条件での揺れ確認」目的。

2026-02-25 - bench.sh 追加（無ウォームアップ・プロセス再起動ベンチ用）

実装:

ルートに bench.sh を追加（実行権限付き）。
target/release/examples/bench_1gb を直接呼び出し、各runを独立プロセスとして実行。
デフォルト条件:
- --size-mib 4096 --runs 5 --iters 1 --warmups 0 --chunk-mib 4 --gpu-subchunk-kib 512 --gpu-fraction 1.0 --mode ratio
機能:
- --mode speed,balanced,ratio の複数モード一括実行
- ログ保存: bench_results/bench_<mode>_<timestamp>.log
- 各モードで mean/median/min/max の要約出力
- デフォルトで COZIP_PROFILE_TIMING/DEEP を解除（--keep-profile-vars で保持可能）

確認:

bash -n bench.sh 通過
./bench.sh --help 表示確認

2026-02-25 - ベンチ結果更新（ユーザー実機 / bench.sh, ratio）

上書き方針:

直前の「ベースライン再計測（ウォームアップなし / プロセス再起動）」は、GPU権限制約付き環境での値だったため、速度比較の基準としては本節のユーザー実機結果を優先する。

入力（ユーザー共有の summary）:

mode=ratio
runs=5
size_mib=4096, iters=1, warmups=0, chunk_mib=4, gpu_subchunk_kib=512, gpu_fraction=1.0

結果:

speedup_compress_x: n=5 mean=1.327 median=1.301 min=1.258 max=1.415
speedup_decompress_x: n=5 mean=1.327 median=1.301 min=1.258 max=1.415
cpu_only_avg_comp_ms: n=5 mean=8295.713 median=8341.937 min=8083.128 max=8367.391
cpu_gpu_avg_comp_ms: n=5 mean=6616.473 median=6590.035 min=6549.101 max=6740.110
gpu_chunks: n=5 mean=283.600 median=285.000 min=268 max=302

運用ルール（今後）:

速度ベンチマークは、原則としてユーザー実機で ./bench.sh を実行して取得した結果を採用する。
評価比較時は runs>=5, iters=1, warmups=0 を基本条件にする（zip実運用前提）。

2026-02-25 - 今後の最適化候補（優先順）

GPU phase1の1パス化（最優先）

tokenize -> finalize -> freq の複数段をできるだけ統合し、グローバルメモリ往復を削減する。
ワークグループ内ヒストグラム（shared memory）で集計し、最後だけglobalに反映する。
期待効果（目安）: 圧縮で +10〜25%。

token_finalize の並列prefix-scan化

直列/疑似直列部分を subgroup + workgroup scan に置換する。
期待効果（目安）: 圧縮で +5〜15%。

GPUバッチの二重バッファ重畳（submit/collect分離）

batch N の待ち中に batch N+1 を投入し、CPU側処理も並行させる。
t_freq_poll_wait_ms を全体時間に埋め込む設計にする。
期待効果（目安）: 圧縮で +5〜15%。

2026-02-25 - 1. GPU phase1最適化に着手（fused phase1パス）

目的:

dynamic経路の tokenize -> token_finalize -> freq のうち、phase1を統合してGPUパスの往復を削減する。

今回の実装（着手版）:

新規 phase1_fused compute pipeline を追加。
各セグメント（TOKEN_FINALIZE_SEGMENT_SIZE）を1 workgroupが担当し、同一パス内で以下を実行:
1. token候補生成（旧tokenize相当）
2. lazy判定付き確定（旧token_finalize相当）
3. litlen/dist頻度集計（workgroupローカル集計後にglobalへ反映）
dynamic圧縮経路 deflate_dynamic_hybrid_batch() では、phase1を
- 旧: tokenize pass + finalize pass + freq pass
- 新: phase1_fused pass に置換。
旧 tokenize/token_finalize/freq pipeline は deep probe 用の互換経路として維持。

補足:

今回は「1.の着手」として、phase1統合を先に導入。
期待していた「完全1パス化」の方向に沿うが、検証容易性のため既存probe資産は保持している。

確認:

cargo check -p cozip_deflate 通過
cargo test -p cozip_deflate -- --nocapture 通過（integration含む）
cargo run --release -p cozip_deflate --example bench_1gb -- --size-mib 64 --iters 1 --warmups 0 --chunk-mib 4 --gpu-subchunk-kib 512 --mode ratio --gpu-fraction 1.0 実行確認（roundtrip成立）

2026-02-25 - 2. token_finalize 並列化（phase1内）

目的:

phase1 fused パス内で、直列 finalize 部分を並列化してGPU実行時間を短縮する。

実装:

phase1 fused shader の finalize を以下に変更:
1. 各laneで lazy判定（mode_lazy_delta）を並列実行し、弱い候補を事前に無効化。
2. 「所有権クレーム」方式で overlap 解決を並列化:
  - token_flags を一時的に owner バッファとして使用。
  - 候補match [i, i+len) に対して atomicMin(owner[p], i) で最左startを確定。
3. owner結果から最終 token を並列確定:
  - owner == 0xFFFFFFFF: literal
  - owner == i: match start
  - それ以外: covered byte（非start）
4. litlen/dist頻度は workgroup ローカル atomic histogram に並列集計し、最後にglobalへ集約。

設計上の注意:

max_storage_buffers_per_shader_stage=8 制約回避のため、追加ownerバッファは導入せず token_flags を再利用。
finalize後に token_flags は 0/1 へ再設定するため後段互換は維持。

確認:

cargo check -p cozip_deflate 通過
cargo test -p cozip_deflate -- --nocapture 通過

追記（同日）:

ユーザー実機ベンチで gpu_chunks=0 が継続し、compressが悪化（speedup_compress_x < 1.0）。
原因は本実装の並列 finalize（ownerクレーム方式）により、ratio検証でGPU結果が不一致となり CPU フォールバックへ落ちること。
対応として、phase1 fused 内の finalize は一旦「直列確定ロジック」に戻した（1.のfused phase1のみ維持）。
以後 2. は「再設計が必要な未完了項目」として扱う。

2026-02-25 - 2. token_finalize 再設計（prefix-scan系の状態伝播）

目的:

ownerクレーム方式の不一致を避けつつ、finalizeを並列化する。

実装（再設計版）:

phase1 fused shader 内 finalize を「レーン分割 + skip状態伝播」へ変更。
1. lazy判定を全laneで並列に事前適用（token_len/token_dist を無効化）。
2. セグメントを 128 lane に等分し、各laneが「入力 skip 値 -> 出力 skip 値」の遷移テーブル（幅33）を作成。
3. lane間の入力 skip を先頭laneから伝播（scan相当のcarry計算）し、各laneの初期状態を確定。
4. 各laneが自区間を独立に finalize（match/literal確定）して token_flags/kind/len/dist を確定。
5. litlen/dist頻度は workgroup local atomic histogram に並列加算し、最後にglobalへ反映。

補足:

token_flags は phase1 fused shader 内のみ array<atomic<u32>> として扱い、atomicStore で 0/1 を確定。
既存bind group構成（StorageBuffer本数）を増やさない形で実装。

確認:

cargo check -p cozip_deflate 通過
cargo test -p cozip_deflate -- --nocapture 通過

2026-02-25 - bench.sh speedup集計バグ修正

事象:

bench.sh の summary で speedup_compress_x と speedup_decompress_x が同値になることがあった。

原因:

sed 抽出式の .* が貪欲で、compress= 側抽出時に後ろの decompress= 側の値を拾っていた。

修正:

bench.sh の speedup抽出を以下へ変更:
- compress: s/.* compress=([0-9.]+)x decompress=.*/\1/
- decompress: s/.* decompress=([0-9.]+)x.*/\1/

補足:

これにより compress / decompress の集計値は独立して算出される。

2026-02-25 - 2.差し戻し後に3. submit/collect重畳を再着手

目的:

2.（token_finalize 並列化）を戻した状態で、3.（submit/collect 分離の重畳）を進める。

実装:

deflate_dynamic_hybrid_batch の freq 段で VecDeque を使った pending 管理を維持。
submit 側で pending が「2wave以上」溜まった時点で、古い 1wave をその場で collect+plan。
重畳 collect 前に device.poll(Maintain::Wait) を入れ、map_async 完了待ちを明示。
prepared の寿命を freq ループ全体に拡張し、未初期化参照バグを修正。

確認:

cargo check -p cozip_deflate 通過。
cargo test -p cozip_deflate -- --nocapture 通過。

2026-02-25 - gpu-dynamic ログ深掘り用の追加計測

目的:

t_freq_poll_wait_ms の支配要因を「待機の種類」まで分解して、次の in-flight 深掘り設計の根拠を作る。

追加した計測（[cozip][timing][gpu-dynamic]）:

in-flight 深さ:
- pending_avg_chunks
- pending_max_chunks
submit→collect 遅延:
- submit_collect_avg_ms
- submit_collect_max_ms
recv 待機内訳:
- recv_immediate
- recv_blocked
- recv_blocked_ms

実装メモ:

PendingDynFreqReadback に submitted_at を保持。
collect 時に submitted_at から遅延を集計。
map_async の待機は try_recv で即時/待機を判別して集計。

確認:

cargo check -p cozip_deflate 通過。
cargo test -p cozip_deflate -- --nocapture 通過。

2026-02-25 - 3. submit/collect重畳の強化（freq Waitバリア縮小）

目的:

t_freq_poll_wait_ms の支配を弱めるため、freq 段の collect を非ブロッキング化し、submit中の Maintain::Wait を避ける。

実装:

PendingDynFreqReadback に litlen_ready/dist_ready を追加。
freq collect を「front要素の try_recv で ready 判定し、ready のみ回収」へ変更。
submitループ中:
- device.poll(Maintain::Poll) 後に ready 分だけ collect。
- これまでの「2wave到達時に Maintain::Wait」を撤去。
最終flush:
- まず Poll で ready 回収。
- 進捗がない場合のみ Maintain::Wait を実施して残件を回収（最小限バリア）。

意図:

Wait を submitパスから退避し、GPU完了待ちをより後段/必要時に限定して重畳性を高める。

確認:

cargo check -p cozip_deflate 通過。
cargo test -p cozip_deflate -- --nocapture 通過。

2026-02-25 - 追加ログの分離（timing_detail）

目的:

追加した詳細計測（pending/submit-collect/recv内訳）を通常ログから分離し、通常運用時の影響を最小化する。

実装:

新規環境変数:
- COZIP_PROFILE_TIMING_DETAIL=1 のときだけ詳細計測を有効化。
COZIP_PROFILE_TIMING=1 のみ:
- 旧来の軽量 gpu-dynamic ログ形式を出力（詳細フィールドなし）。
COZIP_PROFILE_TIMING=1 COZIP_PROFILE_TIMING_DETAIL=1:
- 詳細フィールド付き gpu-dynamic ログを出力。
詳細計測向けの時刻記録（submitted_at）や pending 統計更新は detail 有効時のみ実行。

確認:

cargo check -p cozip_deflate 通過。
cargo test -p cozip_deflate -- --nocapture 通過。

2026-02-25 - scheduler常時ログ拡張 + mode別予約タイムアウト

目的:

圧縮ぶれ原因の切り分け用に、暗黙フォールバック/デモートを scheduler ログで常時可視化する。
GPU_RESERVATION_TIMEOUT_MS=3 固定をやめ、modeごとに適切化する。

実装:

scheduler ログに以下を常時追加:
- demoted_reserved_gpu
- validation_fallback_chunks
GPU検証フォールバック発生時（ratio/balancedの roundtrip mismatch）に validation_fallback_chunks を加算。
watchdog の予約解除タイムアウトを mode 別へ変更:
- Speed: 3ms
- Balanced: 12ms
- Ratio: 32ms
変更点:
- gpu_reservation_timeout_ms(mode) 追加
- compress_scheduler_watchdog に HybridOptions と demoteカウンタを渡す
- compress_gpu_worker_adaptive に validation fallback カウンタを渡す

確認:

cargo check -p cozip_deflate 通過。
cargo test -p cozip_deflate -- --nocapture 通過。

2026-02-25 - 1.2 実装（Ratioデモート抑制）

目的:

demoted_reserved_gpu が過剰に増えることで圧縮速度が不安定になる問題を抑制。

実装:

Ratio の予約タイムアウトを 32ms -> 256ms に引き上げ。
watchdog の Ratio デモートを「GPU進捗停滞時のみ」に制限:
- gpu_last_progress_ms を shared で保持。
- GPU worker が batch 処理完了時に進捗時刻を更新。
- watchdog は now - last_progress >= 512ms のときだけ ReservedGpu -> Pending を許可。

期待:

demoted_reserved_gpu の常時大量発生を抑え、run間の圧縮速度ぶれを低減。

確認:

cargo check -p cozip_deflate 通過。
cargo test -p cozip_deflate -- --nocapture 通過。

2026-02-25 - GPU渋滞時デモートの実効化（DemotedCpu state）

目的:

ReservedGpu -> Pending へ落としたタスクをGPU workerが再取得してしまい、 CPUデモートが効かない競合を解消する。

実装:

CompressTaskState::DemotedCpu を追加。
watchdog のデモート遷移を ReservedGpu -> DemotedCpu に変更。
CPU worker の claim 優先順で DemotedCpu を最優先取得。
GPU worker は DemotedCpu を取得しないため、GPU渋滞時にCPUへ確実に逃がせる。

確認:

cargo check -p cozip_deflate 通過。

2026-02-25 - Ratioスケジューラ安定化（GPU予約バックログ上限）

背景:

同一コードでも run により cpu_chunks が 34 まで落ち、gpu_chunks が 990 へ偏る外れ値が発生。
その際 demoted_reserved_gpu も小さく、GPU待ち渋滞（t_freq_poll_wait_ms 優勢）へ再突入していた。

実装:

Ratio モード限定で、ReservedGpu の予約バックログ上限を導入。
- GPU_RATIO_RESERVED_TARGET_CHUNKS = GPU_BATCH_CHUNKS * 16 (256)
- GPU_RATIO_FORCE_DEMOTE_STEP_CHUNKS = GPU_BATCH_CHUNKS * 4 (64)
watchdog の demote 予算を以下で決定:
- 既存の idle_cpu ベース予算
- 予約過多（ReservedGpu > target）時は、CPUの空きに依存せず min(excess, step) を追加予算として強制demote
これにより GPU 予約の抱え込みを抑え、CPU へ継続的に仕事を供給する。

確認:

cargo check -p cozip_deflate 通過。

2026-02-25 - Ratio demote条件の修正（過剰CPU偏りの是正）

背景:

直前修正後に cpu_chunks~860 / gpu_chunks~160 へ偏り、圧縮が大幅悪化。
原因は、Ratioでも idle CPU ベースdemote が常時動き、予約超過解消後も継続して GPU 予約を削っていたこと。

実装:

watchdog の demote_budget 計算を修正。
Ratio では「予約超過 (ReservedGpu > target) のときのみ」demote。
- 予算: min(backlog_excess, GPU_RATIO_FORCE_DEMOTE_STEP_CHUNKS)
Speed/Balanced は従来どおり idle_cpu ベースdemote。

確認:

cargo check -p cozip_deflate 通過。

2026-02-25 - scheduler実効スループット計測の追加

目的:

cpu_chunks/gpu_chunks が妥当でも圧縮速度が落ちる原因を、CPU/GPU実効処理速度で切り分ける。

実装:

SchedulerPerfCounters を追加し、adaptive scheduler 内で以下を集計:
- CPU: cpu_work_chunks, cpu_work_bytes, cpu_work_ns
- GPU: gpu_batches, gpu_work_chunks, gpu_work_bytes, gpu_work_ns
CPU worker: 各チャンク圧縮の実測時間を cpu_work_ns に加算。
GPU worker: 各バッチ deflate_fixed_literals_batch の実測時間を gpu_work_ns に加算。
scheduler timingログに以下を追加:
- cpu_work_chunks, gpu_work_chunks, gpu_batches
- cpu_work_mib_s, gpu_work_mib_s

期待:

遅い回で「GPUがCPUより遅い」のか「CPU側が並列効率を失っている」のかを明確化できる。

確認:

cargo check -p cozip_deflate 通過。

2026-02-25 - GPUコンテキスト初期化時間ログの追加

目的:

1プロセス1回実行での CPU+GPU 圧縮遅延が、scheduler外のGPU初期化由来かを定量確認する。

実装:

shared_gpu_context() で初回 GpuAssist::new() の経過時間を計測し、 COZIP_PROFILE_TIMING=1 時に以下を出力:
- [cozip][timing] gpu_context_init_ms=... gpu_available=...

確認:

cargo check -p cozip_deflate 通過。

2026-02-25 - 初期化コスト削減（未使用run-start GPUパス削除）

背景:

gpu_context_init_ms=2680ms が単発実行の圧縮性能を大きく悪化させていた。
run_start_positions* 系の match/count/prefix/emit パイプラインは現行ベンチ/圧縮経路で未使用。

実装:

GpuAssist から未使用フィールドを削除:
- match_*, count_*, prefix_*, emit_*
GpuAssist::new_async() で上記4系統の BGL/Shader/Pipeline 作成を削除。
run_start_positions() / run_start_positions_batch() を削除。
既存の deflate 経路で必要な scan_blocks/scan_add は維持。

狙い:

初回 shared_gpu_context() の初期化時間 (gpu_context_init_ms) を直接短縮。

確認:

cargo check -p cozip_deflate 通過。

2026-02-25 - GPU初期化をGPU workerへ遅延移譲（同期初期化の除去）

目的:

compress_hybrid() 呼び出しスレッドでの同期 shared_gpu_context() を避け、 scheduler 起動後に GPU worker 側で初期化することで初回 gpu_context_init_ms をCPU処理と重畳する。

実装:

compress_hybrid() から同期GPU初期化を削除。
- gpu_requested = prefer_gpu && gpu_fraction > 0 で判定し、 GPU対象チャンクがある場合はそのまま adaptive scheduler を起動。
adaptive scheduler に gpu_init_state (INITIALIZING/READY/UNAVAILABLE) を追加。
GPU worker を compress_gpu_worker_adaptive_lazy_init に変更。
- スレッド内で shared_gpu_context() を実行して初期化。
- 初期化成功時: gpu_init_state=READY で従来GPU処理へ移行。
- 初期化失敗時: gpu_init_state=UNAVAILABLE としてCPU側に委譲。
watchdog を gpu_init_state aware に変更。
- UNAVAILABLE では ReservedGpu を即時 DemotedCpu へ落とす。
- READY/INITIALIZING では既存のmode別デモート方針（timeout/backlog）を適用。

意図:

初回起動時に「scheduler開始 -> CPU先行処理 -> GPU初期化完了後にGPU投入」という流れを作り、 2.6s級の初期化コストを体感時間に重畳しやすくする。

確認:

cargo check -p cozip_deflate 通過。

2026-02-25 - scheduler詳細時系列ログの追加（切り分け用）

目的:

「本当にGPU初期化がGPU worker側へ逃げているか」「CPU先行処理が走っているか」を時系列で検証する。

実装:

COZIP_PROFILE_TIMING_DETAIL=1 かつ COZIP_PROFILE_TIMING=1 のときのみ、 scheduler-detail ログを出力:
- start (scheduler起動)
- cpu_first_task (CPU worker初回着手)
- gpu_init_start / gpu_init_ready / gpu_init_unavailable
- gpu_first_batch (GPU初回バッチ着手)
- demote_during_init (初期化中にReservedをCPUへ落とした最初のイベント)
すべて1回だけ出力（mark_once）にしてオーバーヘッドを抑制。

確認:

cargo check -p cozip_deflate 通過。

2026-02-25 - Ratio初期予約数の上限化（初期デモート嵐の抑制）

背景:

COZIP_PROFILE_TIMING_DETAIL=1 で、
- gpu_init_start と同時に demote_during_init
- gpu_init_ready が約2.7s後
- demoted_reserved_gpu が 768 を確認。gpu_fraction=1.0 で初期状態がほぼ全件 ReservedGpu となり、初期化待ち中に大量デモートする無駄が発生していた。

実装:

adaptive scheduler の初期state付与を変更。
Ratio では初期 ReservedGpu を GPU_RATIO_RESERVED_TARGET_CHUNKS (256) に制限し、それ以外は最初から Pending に置く。
これにより watchdog の初期デモート量を減らし、CPU先行実行の立ち上がりを安定化。

確認:

cargo check -p cozip_deflate 通過。

2026-02-25 - 初期予約のcold-start制御（warm/coldで分岐）

背景:

initial_reserved=256 + demoted_reserved_gpu=0 でも初回実行は遅く、 cold-startでは予約チャンクを待つ構造が残っていた。

実装:

gpu_context_cached_available() を追加し、GPUコンテキストがプロセス内で既に温まっているか判定。
Ratio の初期予約数を warm/cold で分岐:
- warm (gpu_cached_ready=true): GPU_RATIO_RESERVED_TARGET_CHUNKS (256)
- cold (gpu_cached_ready=false): GPU_BATCH_CHUNKS (16)
scheduler-detail start に gpu_cached_ready を追加。

狙い:

cold-start時の「初期予約待ち」を最小化し、CPU先行実行を阻害しない。

2026-02-25 - cold-start初期予約数チューニング（16->32）

背景:

cold-start initial_reserved=16 は安定化に効いたが、gpu_chunks が下がりすぎて圧縮上振れ余地が残った。

実装:

GPU_COLD_START_RESERVED_CHUNKS = GPU_BATCH_CHUNKS * 2 を追加。
cold-start時の初期予約数を 16 -> 32 に調整。

狙い:

起動直後のCPU先行性を維持しつつ、GPU投入を少し早めて圧縮スループットを改善。

確認:

cargo check -p cozip_deflate 通過。

2026-02-25 - cold-start初期予約の既定値を16へ戻し

背景:

実測上 32 の優位性が環境依存で、低性能GPUでの安全側設定を優先する判断。

実装:

GPU_COLD_START_RESERVED_CHUNKS の既定値を GPU_BATCH_CHUNKS (16) に戻し。
必要時は COZIP_GPU_COLD_RESERVED_CHUNKS で上書き可能な構成は維持。

2026-02-26 - `gpu_chunks` 偏りの主因を特定（GPU初期化時間ジッタ）

背景:

同一コマンド・同一設定でも gpu_chunks が 200台と300台で偏る現象が継続。
scheduler調整のみでは再現パターン（初回低め、後続で高め）が解消しなかった。

観測:

COZIP_PROFILE_TIMING=1 COZIP_PROFILE_TIMING_DETAIL=1 で、同一コマンドにて:
- 遅い回: gpu_context_init_ms=2646.749
- 速い回: gpu_context_init_ms=208.249
この差に対応して gpu_ready_snapshot が大きく変化:
- 遅い回: done=320（GPU ready前にCPUが先行消化）
- 速い回: done=7
最終配分も連動:
- 遅い回: gpu_chunks=240
- 速い回: gpu_chunks=336

結論:

gpu_chunks の主な偏り要因は scheduler 内の量子化より、GpuAssist 初期化完了時刻のジッタ。
要因の中心は wgpu/Vulkan/driver 側の cold/warm 差（パイプライン・ドライバキャッシュ等）である可能性が高い。
実装側はこの外部ジッタを観測しやすい実行形態（bench.sh の runごとプロセス再起動）であるため、偏りが顕在化した。

補足:

解凍ベンチが比較的安定していたのは、初期化ジッタの影響が主に圧縮側スケジューリングへ現れるためと整合する。

2026-02-28 - 独自GDeflateハイブリッド設計（CPU+GPU実処理参加）を追加

決定事項:

docs/gdeflate-hybrid-cpu-gpu-design.md を新規追加。
当面は互換性後回しで、独自形式 CGDF v0 を先行実装する方針。
CPUを管理専用にせず、CPU/GPUが同一キューから実処理タスクを消費する。
スケジューラーは CoZipDeflate の GlobalQueueLocalBuffers 準拠を維持する。

補足:

圧縮・解凍ともに同じ queue/ready/error の共有構造を採用する前提を明記。
後続フェーズでMS互換対応を行う段階計画（G0-G4）を定義。

2026-02-28 - cozip_gdeflate G1着手（Hybrid圧縮の最小実装）

実装:

src/cozip_gdeflate に G1の最小構成を実装。
CGDF v0 は維持したまま、圧縮経路に共有キュー方式を導入。
- CPUワーカー: deflate_compress_chunk を実行
- GPUワーカー: CoZipDeflate を利用したGPU支援圧縮を実行（失敗時はCPUフォールバック）
- Writer: index順に再整列してフレーム確定
GdHybridStats を追加し、CPU/GPUチャンク数・busy/wait・batch統計を収集。

確認:

cargo test -p cozip_gdeflate 通過（6 tests）。
cargo check -p cozip_gdeflate 通過。

2026-02-28 - cozip_gdeflate を純Rust GDeflate CPU encode/decode 初版へ切替

実装:

cozip_deflate 依存を除去し、src/cozip_gdeflate/src/lib.rs を純Rust実装へ更新。
DirectStorage参照実装の TileStream ヘッダ形式（id/magic/numTiles/tile metadata）を採用。
CPU encode/decode を実装（現時点は stored-block ベース）。
- タイル(64KiB)分割
- 32 sub-stream への割当
- タイルストリームの offset テーブル生成
- 復号時の逆変換とサイズ検証

確認:

cargo test -p cozip_gdeflate 通過（4 tests）。

2026-03-07 - PDeflate GPU sparse direct output の設計分解を追加

背景:

speed モードでは GPU table build を無効化済み。
それでも GPU圧縮の主要律速が sparse_lens_wait に残っている。
現行 integrated sparse path は result readback -> payload readback の 2 段 readback で、device.poll(Maintain::Wait) が強い待ちを作っている。

決定:

全面 worst-case 固定長ではなく、サイズクラス別 direct output を本命案とする。
まずは docs/pdeflate-gpu-sparse-direct-output-design.md と docs/tasks/pdeflate-gpu-sparse-direct-output-tasks.md で、段階導入の設計とフェーズを定義した。

狙い:

sparse_lens_wait の根本原因である「長さ待ち後に payload 回収計画を作る」依存を壊す。
VRAM 悪化を避けつつ、2 段 readback を解消する。

2026-03-07 - PDeflate GPU sparse direct output Phase 0/1 を実装

実装:

GpuMatchInput に compression_mode を追加し、legacy 側の GPU圧縮経路へ伝搬。
pdeflate::gpu に sparse payload size class を追加。
- Tiny <= 256KiB
- Small <= 1MiB
- Medium <= 2MiB
- Large <= 4MiB
- XLarge > 4MiB
compute_matches_encode_and_pack_sparse_batch で、chunk ごとに conservative な sparse payload 推定値と class を計算するようにした。
実 payload 回収時に、予測 class / cap class / 実 class のヒストグラムと、under/over estimate 量を集計するようにした。
新しい計測ログを追加:
- sparse-class-breakdown
- gpu-pack-batch への predicted/actual/cap average と under/over 量の追加

狙い:

Phase 2 以降の classed direct output 設計に必要な実分布を取る。
「VRAM を爆発させずに事前 slot 割当できる class 境界」を決める。

確認:

cargo check -p cozip_pdeflate 通過。
cargo build --release --example bench_pdeflate -p cozip_pdeflate 通過。

2026-03-07 - PDeflate GPU sparse direct output Phase 2 を実装

実装:

integrated sparse path の out_buffer 配置を、予測 class ベースの slot layout へ変更。
GPU_SPARSE_PACK_BATCH_DESC_WORDS を拡張し、class/slot 情報を desc へ載せる土台を追加。
chunk ごとの slot_cap_bytes を class 容量から決め、out_base_word と total_bytes_cap を classed slot に合わせて再配置するようにした。
sparse-class-layout ログを追加し、旧レイアウト比の out_total_mib 縮小率を見えるようにした。

注意:

この段階では 2 段 readback 自体はまだ残っている。
したがって sparse_lens_wait の本体削減は Phase 4 以降。
期待値は「まず VRAM と output slot 容量を削る」であり、速度改善は副次的。

確認:

cargo check -p cozip_pdeflate 通過。
cargo build --release --example bench_pdeflate -p cozip_pdeflate 通過。

2026-03-08 CoZip PDeflate directory mode

cozip の PDeflate backend に directory compress/decompress を追加した
実装方式は ZIP 互換ではなく、cozip 内部の streaming archive を先に作ってから PDeflate へ流す archive -> compress 方式
解凍はその逆で、PDeflate を stream 解凍しつつ内部 archive を逐次展開する
これにより一時 spool file なしで directory mode を扱える

2026-03-10 PDeflate single-file progress metadata

PDS0 ストリームヘッダに、任意の uncompressed_size: u64 metadata を追加した
単一ファイル PDeflate 圧縮では入力ファイルサイズをこの metadata として埋め込む
単一ファイル PDeflate 解凍では、この metadata が存在する場合に事前全走査なしで write 基準の進捗率を出せる
metadata が存在しない旧ストリームは引き続き解凍可能で、進捗率は不定扱いとする

2026-03-11 PDeflate single-file mmap parallel write

単一ファイル PDeflate 解凍に decompress_file_parallel_write 系 API を追加した
出力先ファイルを memmap2 で map し、各チャンクの chunk_uncompressed_len 累積から算出した出力 offset へ直接並列書き込みする
解凍 worker と書き込み worker を分離し、復元済みチャンクは write queue 経由で SSD へ並列反映する
書き込み queue backlog の上限は固定 2 GiB とし、Desktop 側へ backlog 警告を出せるようにした
cozip_desktop の単一ファイル PDeflate 解凍はこの経路を優先使用し、書き込みスレッド数は PDeflateOptions.parallel_write_threads で調整できる

2026-03-11 cozip_util ParallelFileWriter

src/cozip_util クレートを追加し、ParallelFileWriter を OS ごとの positional write backend として切り出した
Windows は seek_write、Unix は write_all_at、その他 OS は mutex + seek/write のフォールバック実装
PDeflate 単一ファイル解凍の parallel write は memmap2 直書きから cozip_util::ParallelFileWriter ベースへ置き換えた

2026-03-11 cozip_util async file backends

cozip_util::ParallelFileWriter の backend を submit/completion 型に整理した
Windows backend は ReOpenFile(..., FILE_FLAG_OVERLAPPED) + CreateIoCompletionPort + WriteFile(OVERLAPPED) に変更した
Linux backend は io_uring crate を使う専用 thread backend を追加した
非 Windows / 非 Linux は従来どおり mutex + positional write のフォールバックを維持する

2026-03-11 PDeflate single-file filename metadata

PDeflate 単一ファイル stream header に任意の UTF-8 file_name metadata を追加した
単一ファイル圧縮では元ファイル名を埋め込み、解凍時はその名前を優先して復元する
旧ストリームで metadata がない場合は、従来どおりアーカイブ名 stem から復元名を推定する

FilesExpand file tree

context-log.md

Latest commit

History

context-log.md

File metadata and controls

Context Log

2026-02-24 - 初回設計ドラフト作成

決定事項

採用理由

トレードオフ

次アクション

2026-02-24 - workspace分割 + Deflate実装着手

決定事項

補足

次アクション

2026-02-24 - CPU vs CPU+GPU 統合テスト追加

決定事項

仕様

実行コマンド

2026-02-24 - GPU実タスク設計(大チャンク転送 + GPU内細分化)

決定事項

採用理由

追記したドキュメント

次アクション

2026-02-24 - CPU全力 + GPU全力 実装(第1段)

決定事項

GPU実タスク(圧縮/解凍)

互換性

検証

2026-02-24 - 速度ベンチマーク実施(Release)

実行条件

計測結果

所見

2026-02-24 - 1GB比較用ベンチ追加

決定事項

実行コマンド

次アクション

2026-02-24 - 1GBベンチのデフォルト実行回数を短縮

決定事項

理由

補足

2026-02-24 - GPU Deflate本体設計(独立チャンク実行 + 連結)

決定事項

追記したドキュメント

次アクション

2026-02-24 - GPU Deflate本体実装(第2段)

実装内容

主要ポイント

検証

2026-02-27 - D0実装（CZDIメタデータ基盤）

実装内容

ZIP統合

補足

補足

2026-02-24 - 改善1/3実装 (GPUトークン品質 + 大きめチャンク)

実装内容

計測(bench_hybrid, release, GPU実機)

所見

追記テンプレート

2026-02-24 - 改善2/3着手 (GPUバッチsubmit + 転送前処理削減)

実装内容

検証

メモ

2026-02-24 - GPU Deflate完結経路(固定Huffman literal) + readback最小化

実装内容

検証

直近ベンチ(bench_hybrid)

メモ

2026-02-24 - 追加改善(継続): 安全側の高速化調整

対応

重要メモ

検証

2026-02-24 - GPU len/dist 実装 (RLE系 match token 追加)

実装内容

バグ修正

検証

観測

2026-02-24 - GPU len/dist 強化 (candidate + finalize 2段)

実装内容

2026-02-24 - CPU全力 + GPU全力実装(第1段)