Commit Graph

551 Commits

Author SHA1 Message Date
root
89f1e5b743 build: trigger rebuild after merge reconciliation 2026-08-13 10:08:16 +00:00
Claude
61bfb8df29 fix: wire CCCL preload into build+launch chain + pre-submission verification
- patch_ops.sh: call cccl_preload/build_cccl_preload.sh (new CCCL deps)
  instead of old build_cccl_preload_allocator.sh (mock)
- computility-run.yaml: add LD_PRELOAD + CCCL_ALLOC_DISABLE env vars
- Remove old mock files: cccl_preload_allocator.cu, build script, test
- .dockerignore: exclude cccl_upstream/ upstream_ref/ vllm/ *.zip
- verify_submission.sh: 31-point pre-submission check
  (file structure, CCCL chain, path matching, prebuilt integrity,
   corex imports, docker context, GPU smoke test)
2026-08-13 10:08:16 +00:00
root
3c49308449 Merge branch 'main' of https://dev.modelhub.org.cn/dylanyunlong/project_6 2026-08-13 09:54:10 +00:00
dylanyunlon
a6b5891bfc feat: CCCL CachingDeviceAllocator preload — 完整依赖链 288 files
从 cccl_upstream 递归追踪 cub/util_allocator.cuh 的全部 include 依赖:
  cub/         9 files (config, util_*, version, detect_cuda_runtime)
  cuda/        libcudacxx type_traits, concepts, algorithm, iterator...
  nv/          target macros, preprocessor

总计 288 个头文件 (1.4MB),打包到 include/ 目录,编译时 -I include
即可完全脱离 CCCL 原始目录结构。

.cu 文件直接 #include <cub/util_allocator.cuh>,
走原版 CUB CachingDeviceAllocator,零 mock。

BI-V100 参数: growth=2 bins=[8..32] max_cached=8GB/device
2026-08-13 09:53:42 +00:00
dylanyunlon
8dc6462a2b feat: CCCL CachingDeviceAllocator LD_PRELOAD — bypass CoreX expandable_segments ASSERT
从 CCCL upstream cub/cub/util_allocator.cuh 提取 CachingDeviceAllocator
核心算法,去掉所有 CUB/CCCL 宏依赖,编译为独立 .so。

用 LD_PRELOAD 拦截 cudaMalloc/cudaFree,路由到 CUB 的 geometric-bin
缓存分配器。同时在 constructor 中 strip PYTORCH_CUDA_ALLOC_CONF 里的
expandable_segments 配置,避免 CoreX CUDACachingAllocator.cpp:545 ASSERT。

BI-V100 调优参数:
  bin_growth=8, min_bin=3 (512B), max_bin=13 (~550MB)
  max_cached_bytes=4GB per device (32GB卡的合理上限)

真机测试步骤:
  1. bash build_cccl_preload.sh
  2. LD_PRELOAD=./libcccl_allocator.so CCCL_ALLOC_DEBUG=1 \
     PYTORCH_CUDA_ALLOC_CONF=expandable_segments:True \
     python3 verify_preload.py
2026-08-13 09:53:42 +00:00
dylanyunlon
887e0981ad fix: remove expandable_segments — CoreX CUDACachingAllocator不支持
BI-V100 CoreX PyTorch的CUDACachingAllocator.cpp:545没有实现
expandable segment特性,导致模型加载阶段(VocabParallelEmbedding)
直接INTERNAL ASSERT FAILED崩溃。

替换为max_split_size_mb:512减少内存碎片化。
2026-08-13 09:53:37 +00:00
Claude
9cfc6c72c0 feat(CCCL): LD_PRELOAD CachingDeviceAllocator — intercept cudaMalloc/cudaFree
Route C: replace PyTorch's cudaMalloc/cudaFree with CCCL CUB's
CachingDeviceAllocator via LD_PRELOAD. Eliminates driver-level allocation
overhead by reusing freed GPU memory from a bin-based cache.

Based on cccl_upstream/cub/cub/util_allocator.cuh (901 lines).
Self-contained .so with no CCCL header dependencies at compile time.

Files:
- cccl_preload_allocator.cu: the allocator (405 lines)
- build_cccl_preload_allocator.sh: build script (corex clang++ or g++ fallback)
- test_cccl_preload.sh: smoke test suite for BI-V100
- patch_ops.sh: build during docker build
- computility-run.yaml: LD_PRELOAD env var for runtime

Config via env:
  CCCL_ALLOC_BIN_GROWTH=8, MIN_BIN=3, MAX_BIN=13, MAX_CACHED_MB=4096

Test on real machine:
  cd qwen3_6_scripts && bash test_cccl_preload.sh
2026-08-13 09:53:37 +00:00
root
ca42633148 build: trigger rebuild after merge reconciliation 2026-08-13 09:53:19 +00:00
root
7e4e04b7c6 build: trigger rebuild after merge reconciliation 2026-08-13 09:53:19 +00:00
dylanyunlon
089e810984 feat: CCCL CachingDeviceAllocator preload — 完整依赖链 288 files
从 cccl_upstream 递归追踪 cub/util_allocator.cuh 的全部 include 依赖:
  cub/         9 files (config, util_*, version, detect_cuda_runtime)
  cuda/        libcudacxx type_traits, concepts, algorithm, iterator...
  nv/          target macros, preprocessor

总计 288 个头文件 (1.4MB),打包到 include/ 目录,编译时 -I include
即可完全脱离 CCCL 原始目录结构。

.cu 文件直接 #include <cub/util_allocator.cuh>,
走原版 CUB CachingDeviceAllocator,零 mock。

BI-V100 参数: growth=2 bins=[8..32] max_cached=8GB/device
2026-08-13 09:53:19 +00:00
dylanyunlon
e7c703ef94 feat: CCCL CachingDeviceAllocator LD_PRELOAD — bypass CoreX expandable_segments ASSERT
从 CCCL upstream cub/cub/util_allocator.cuh 提取 CachingDeviceAllocator
核心算法,去掉所有 CUB/CCCL 宏依赖,编译为独立 .so。

用 LD_PRELOAD 拦截 cudaMalloc/cudaFree,路由到 CUB 的 geometric-bin
缓存分配器。同时在 constructor 中 strip PYTORCH_CUDA_ALLOC_CONF 里的
expandable_segments 配置,避免 CoreX CUDACachingAllocator.cpp:545 ASSERT。

BI-V100 调优参数:
  bin_growth=8, min_bin=3 (512B), max_bin=13 (~550MB)
  max_cached_bytes=4GB per device (32GB卡的合理上限)

真机测试步骤:
  1. bash build_cccl_preload.sh
  2. LD_PRELOAD=./libcccl_allocator.so CCCL_ALLOC_DEBUG=1 \
     PYTORCH_CUDA_ALLOC_CONF=expandable_segments:True \
     python3 verify_preload.py
2026-08-13 09:53:19 +00:00
dylanyunlon
c1e7065076 fix: remove expandable_segments — CoreX CUDACachingAllocator不支持
BI-V100 CoreX PyTorch的CUDACachingAllocator.cpp:545没有实现
expandable segment特性,导致模型加载阶段(VocabParallelEmbedding)
直接INTERNAL ASSERT FAILED崩溃。

替换为max_split_size_mb:512减少内存碎片化。
2026-08-13 09:53:01 +00:00
Claude
1ea2100cb8 feat(CCCL): LD_PRELOAD CachingDeviceAllocator — intercept cudaMalloc/cudaFree
Route C: replace PyTorch's cudaMalloc/cudaFree with CCCL CUB's
CachingDeviceAllocator via LD_PRELOAD. Eliminates driver-level allocation
overhead by reusing freed GPU memory from a bin-based cache.

Based on cccl_upstream/cub/cub/util_allocator.cuh (901 lines).
Self-contained .so with no CCCL header dependencies at compile time.

Files:
- cccl_preload_allocator.cu: the allocator (405 lines)
- build_cccl_preload_allocator.sh: build script (corex clang++ or g++ fallback)
- test_cccl_preload.sh: smoke test suite for BI-V100
- patch_ops.sh: build during docker build
- computility-run.yaml: LD_PRELOAD env var for runtime

Config via env:
  CCCL_ALLOC_BIN_GROWTH=8, MIN_BIN=3, MAX_BIN=13, MAX_CACHED_MB=4096

Test on real machine:
  cd qwen3_6_scripts && bash test_cccl_preload.sh
2026-08-13 09:53:01 +00:00
root
abbc13c4d5 build: trigger rebuild after merge reconciliation 2026-08-13 09:53:01 +00:00
root
a8304bf906 build: trigger rebuild after merge reconciliation 2026-08-13 09:53:01 +00:00
dylanyunlon
ddfd24da27 fix: sync to real-machine verified version — ALL TESTS PASSED
真机验证通过的精确版本:
- CUB_NS_QUALIFIER (不是 cub::)
- thread_local inside_cub reentrant guard
- 去掉 -D_CCCL_COMPILER_GCC=1
- total_mem → total_memory

BI-V100 32GB × Iluvatar, CoreX clang++ 编译 51864 bytes .so
expandable_segments:True 被 strip, CUB allocator 接管, 缓存复用确认
2026-08-13 09:52:15 +00:00
root
ffe764f8b2 Merge branch 'main' of https://dev.modelhub.org.cn/dylanyunlong/project_6 2026-08-13 09:46:52 +00:00
dylanyunlon
2cbae09d5e feat: CCCL CachingDeviceAllocator preload — 完整依赖链 288 files
从 cccl_upstream 递归追踪 cub/util_allocator.cuh 的全部 include 依赖:
  cub/         9 files (config, util_*, version, detect_cuda_runtime)
  cuda/        libcudacxx type_traits, concepts, algorithm, iterator...
  nv/          target macros, preprocessor

总计 288 个头文件 (1.4MB),打包到 include/ 目录,编译时 -I include
即可完全脱离 CCCL 原始目录结构。

.cu 文件直接 #include <cub/util_allocator.cuh>,
走原版 CUB CachingDeviceAllocator,零 mock。

BI-V100 参数: growth=2 bins=[8..32] max_cached=8GB/device
2026-08-13 09:45:45 +00:00
dylanyunlon
5449456534 feat: CCCL CachingDeviceAllocator LD_PRELOAD — bypass CoreX expandable_segments ASSERT
从 CCCL upstream cub/cub/util_allocator.cuh 提取 CachingDeviceAllocator
核心算法,去掉所有 CUB/CCCL 宏依赖,编译为独立 .so。

用 LD_PRELOAD 拦截 cudaMalloc/cudaFree,路由到 CUB 的 geometric-bin
缓存分配器。同时在 constructor 中 strip PYTORCH_CUDA_ALLOC_CONF 里的
expandable_segments 配置,避免 CoreX CUDACachingAllocator.cpp:545 ASSERT。

BI-V100 调优参数:
  bin_growth=8, min_bin=3 (512B), max_bin=13 (~550MB)
  max_cached_bytes=4GB per device (32GB卡的合理上限)

真机测试步骤:
  1. bash build_cccl_preload.sh
  2. LD_PRELOAD=./libcccl_allocator.so CCCL_ALLOC_DEBUG=1 \
     PYTORCH_CUDA_ALLOC_CONF=expandable_segments:True \
     python3 verify_preload.py
2026-08-13 09:45:45 +00:00
dylanyunlon
1e8d1ddfa3 fix: remove expandable_segments — CoreX CUDACachingAllocator不支持
BI-V100 CoreX PyTorch的CUDACachingAllocator.cpp:545没有实现
expandable segment特性,导致模型加载阶段(VocabParallelEmbedding)
直接INTERNAL ASSERT FAILED崩溃。

替换为max_split_size_mb:512减少内存碎片化。
2026-08-13 09:43:35 +00:00
Claude
bf0ad625c3 feat(CCCL): LD_PRELOAD CachingDeviceAllocator — intercept cudaMalloc/cudaFree
Route C: replace PyTorch's cudaMalloc/cudaFree with CCCL CUB's
CachingDeviceAllocator via LD_PRELOAD. Eliminates driver-level allocation
overhead by reusing freed GPU memory from a bin-based cache.

Based on cccl_upstream/cub/cub/util_allocator.cuh (901 lines).
Self-contained .so with no CCCL header dependencies at compile time.

Files:
- cccl_preload_allocator.cu: the allocator (405 lines)
- build_cccl_preload_allocator.sh: build script (corex clang++ or g++ fallback)
- test_cccl_preload.sh: smoke test suite for BI-V100
- patch_ops.sh: build during docker build
- computility-run.yaml: LD_PRELOAD env var for runtime

Config via env:
  CCCL_ALLOC_BIN_GROWTH=8, MIN_BIN=3, MAX_BIN=13, MAX_CACHED_MB=4096

Test on real machine:
  cd qwen3_6_scripts && bash test_cccl_preload.sh
2026-08-13 09:43:35 +00:00
root
1aa6e16d4e build: trigger rebuild after merge reconciliation 2026-08-13 09:43:18 +00:00
root
93fb9bd8b7 build: trigger rebuild after merge reconciliation 2026-08-13 09:43:18 +00:00
dylanyunlon
93e498197a fix: thread_local reentrant guard — prevent cudaMalloc infinite recursion
CUB CachingDeviceAllocator 内部在 cache miss 时调 cudaMalloc,
被我们的 LD_PRELOAD 再次拦截 → DeviceAllocate → cudaMalloc → 无限递归 → segfault。

加 thread_local bool inside_cub 标志:
  外部调用 → CUB allocator (带缓存)
  CUB 内部调用 → 直接走 dlsym(RTLD_NEXT) 的真实 cudaMalloc
2026-08-13 09:42:26 +00:00
root
19946d4c72 merge: reconcile modelhub with github (keep CUB_NS_QUALIFIER fix) 2026-08-13 09:35:35 +00:00
root
188d3654ef build: trigger rebuild after merge reconciliation 2026-08-13 09:34:49 +00:00
dylanyunlon
0ac118911d fix: CUB_NS_QUALIFIER for wrapped namespace + drop _CCCL_COMPILER_GCC
CoreX clang++ 不是 GCC,-D_CCCL_COMPILER_GCC=1 和 CCCL 自己的
compiler detection 冲突。

CUB_WRAPPED_NAMESPACE=cccl_preload 使得命名空间变成 cccl_preload::cub,
用 CUB_NS_QUALIFIER 宏自动解析正确的命名空间。
2026-08-13 09:31:45 +00:00
dylanyunlon
e1cdbfc167 feat: CCCL CachingDeviceAllocator preload — 完整依赖链 288 files
从 cccl_upstream 递归追踪 cub/util_allocator.cuh 的全部 include 依赖:
  cub/         9 files (config, util_*, version, detect_cuda_runtime)
  cuda/        libcudacxx type_traits, concepts, algorithm, iterator...
  nv/          target macros, preprocessor

总计 288 个头文件 (1.4MB),打包到 include/ 目录,编译时 -I include
即可完全脱离 CCCL 原始目录结构。

.cu 文件直接 #include <cub/util_allocator.cuh>,
走原版 CUB CachingDeviceAllocator,零 mock。

BI-V100 参数: growth=2 bins=[8..32] max_cached=8GB/device
2026-08-13 09:30:44 +00:00
dylanyunlon
8eba8df925 feat: CCCL CachingDeviceAllocator LD_PRELOAD — bypass CoreX expandable_segments ASSERT
从 CCCL upstream cub/cub/util_allocator.cuh 提取 CachingDeviceAllocator
核心算法,去掉所有 CUB/CCCL 宏依赖,编译为独立 .so。

用 LD_PRELOAD 拦截 cudaMalloc/cudaFree,路由到 CUB 的 geometric-bin
缓存分配器。同时在 constructor 中 strip PYTORCH_CUDA_ALLOC_CONF 里的
expandable_segments 配置,避免 CoreX CUDACachingAllocator.cpp:545 ASSERT。

BI-V100 调优参数:
  bin_growth=8, min_bin=3 (512B), max_bin=13 (~550MB)
  max_cached_bytes=4GB per device (32GB卡的合理上限)

真机测试步骤:
  1. bash build_cccl_preload.sh
  2. LD_PRELOAD=./libcccl_allocator.so CCCL_ALLOC_DEBUG=1 \
     PYTORCH_CUDA_ALLOC_CONF=expandable_segments:True \
     python3 verify_preload.py
2026-08-13 09:30:43 +00:00
dylanyunlon
1a8d086c4a fix: remove expandable_segments — CoreX CUDACachingAllocator不支持
BI-V100 CoreX PyTorch的CUDACachingAllocator.cpp:545没有实现
expandable segment特性,导致模型加载阶段(VocabParallelEmbedding)
直接INTERNAL ASSERT FAILED崩溃。

替换为max_split_size_mb:512减少内存碎片化。
2026-08-13 09:30:43 +00:00
Claude
d036646726 feat(CCCL): LD_PRELOAD CachingDeviceAllocator — intercept cudaMalloc/cudaFree
Route C: replace PyTorch's cudaMalloc/cudaFree with CCCL CUB's
CachingDeviceAllocator via LD_PRELOAD. Eliminates driver-level allocation
overhead by reusing freed GPU memory from a bin-based cache.

Based on cccl_upstream/cub/cub/util_allocator.cuh (901 lines).
Self-contained .so with no CCCL header dependencies at compile time.

Files:
- cccl_preload_allocator.cu: the allocator (405 lines)
- build_cccl_preload_allocator.sh: build script (corex clang++ or g++ fallback)
- test_cccl_preload.sh: smoke test suite for BI-V100
- patch_ops.sh: build during docker build
- computility-run.yaml: LD_PRELOAD env var for runtime

Config via env:
  CCCL_ALLOC_BIN_GROWTH=8, MIN_BIN=3, MAX_BIN=13, MAX_CACHED_MB=4096

Test on real machine:
  cd qwen3_6_scripts && bash test_cccl_preload.sh
2026-08-13 09:30:43 +00:00
dylanyunlon
8d6f9eaeb0 feat: CCCL CachingDeviceAllocator preload — 完整依赖链 288 files
从 cccl_upstream 递归追踪 cub/util_allocator.cuh 的全部 include 依赖:
  cub/         9 files (config, util_*, version, detect_cuda_runtime)
  cuda/        libcudacxx type_traits, concepts, algorithm, iterator...
  nv/          target macros, preprocessor

总计 288 个头文件 (1.4MB),打包到 include/ 目录,编译时 -I include
即可完全脱离 CCCL 原始目录结构。

.cu 文件直接 #include <cub/util_allocator.cuh>,
走原版 CUB CachingDeviceAllocator,零 mock。

BI-V100 参数: growth=2 bins=[8..32] max_cached=8GB/device
2026-08-13 09:26:41 +00:00
dylanyunlon
967d572073 feat: CCCL CachingDeviceAllocator LD_PRELOAD — bypass CoreX expandable_segments ASSERT
从 CCCL upstream cub/cub/util_allocator.cuh 提取 CachingDeviceAllocator
核心算法,去掉所有 CUB/CCCL 宏依赖,编译为独立 .so。

用 LD_PRELOAD 拦截 cudaMalloc/cudaFree,路由到 CUB 的 geometric-bin
缓存分配器。同时在 constructor 中 strip PYTORCH_CUDA_ALLOC_CONF 里的
expandable_segments 配置,避免 CoreX CUDACachingAllocator.cpp:545 ASSERT。

BI-V100 调优参数:
  bin_growth=8, min_bin=3 (512B), max_bin=13 (~550MB)
  max_cached_bytes=4GB per device (32GB卡的合理上限)

真机测试步骤:
  1. bash build_cccl_preload.sh
  2. LD_PRELOAD=./libcccl_allocator.so CCCL_ALLOC_DEBUG=1 \
     PYTORCH_CUDA_ALLOC_CONF=expandable_segments:True \
     python3 verify_preload.py
2026-08-13 09:26:41 +00:00
dylanyunlon
502ea2fc96 fix: remove expandable_segments — CoreX CUDACachingAllocator不支持
BI-V100 CoreX PyTorch的CUDACachingAllocator.cpp:545没有实现
expandable segment特性,导致模型加载阶段(VocabParallelEmbedding)
直接INTERNAL ASSERT FAILED崩溃。

替换为max_split_size_mb:512减少内存碎片化。
2026-08-13 09:26:41 +00:00
Claude
327c2c9044 feat(CCCL): LD_PRELOAD CachingDeviceAllocator — intercept cudaMalloc/cudaFree
Route C: replace PyTorch's cudaMalloc/cudaFree with CCCL CUB's
CachingDeviceAllocator via LD_PRELOAD. Eliminates driver-level allocation
overhead by reusing freed GPU memory from a bin-based cache.

Based on cccl_upstream/cub/cub/util_allocator.cuh (901 lines).
Self-contained .so with no CCCL header dependencies at compile time.

Files:
- cccl_preload_allocator.cu: the allocator (405 lines)
- build_cccl_preload_allocator.sh: build script (corex clang++ or g++ fallback)
- test_cccl_preload.sh: smoke test suite for BI-V100
- patch_ops.sh: build during docker build
- computility-run.yaml: LD_PRELOAD env var for runtime

Config via env:
  CCCL_ALLOC_BIN_GROWTH=8, MIN_BIN=3, MAX_BIN=13, MAX_CACHED_MB=4096

Test on real machine:
  cd qwen3_6_scripts && bash test_cccl_preload.sh
2026-08-13 09:21:52 +00:00
root
05713aa7cb build: trigger rebuild after merge reconciliation 2026-08-13 08:46:38 +00:00
project6-dev
a1ae6e366f merge: reconcile squashed commit with modelhub history 2026-08-13 07:10:03 +00:00
project6-dev
d2b4df54ff perf: native ixformer decode — v1 ≤32K, v2 >32K (no Python fallback)
Decode path:
- ≤32K: paged_attention_v1 (5D KV layout, x=8) — verified on real BI-V100
- >32K: paged_attention_v2 (5D→4D permute) — verified 65K on real BI-V100
- Removes _forward_decode_pytorch Python fallback entirely

Verified: v1 passes ctx=1024..32768, v2 passes ctx=32768..65536+
flash_attn_varlen_func prefill already merged in prior commit (ad6863ed).
2026-08-13 07:09:50 +00:00
project6-dev
f28223c9da perf: native ixformer decode (v1 ≤32K, v2 >32K) + flash_attn_varlen prefill
Replaces all Python PyTorch fallback attention with native ixformer kernels:

Decode path:
- ≤32K: paged_attention_v1 (5D KV layout, x=8) — verified on real BI-V100
- >32K: paged_attention_v2 (5D→4D permute) — verified 65K+ on real BI-V100
- Removes _forward_decode_pytorch Python fallback entirely

Prefill path (profiling):
- _run_sdpa_fallback now uses ixformer.flash_attn_varlen_func
- head_dim=256 verified correct (diff<0.004) and 1.7x faster than PyTorch
- Falls back to Q-tiling pure-math if ixformer unavailable

Also includes: MoE kernel integration, GDN C++ kernels, diagnostic scripts,
xllm upstream layer/kernel references, .dockerignore cleanup.

All changes verified on real BI-V100 hardware (single card).
2026-08-13 07:04:21 +00:00
Claude
e78fa560c8 feat: wire corex_gdn_chunk_recurrent C++ kernel into GDN prefill path
- patch_ops.sh: build corex_gdn_chunk_recurrent.so alongside moe_index_combine
- qwen3_5.py: import corex_gdn_chunk_recurrent, use C++ version for prefill
  chunks instead of Python _torch_chunk_gated_delta_rule
- C++ version from xllm upstream avoids Python loop overhead and has proper
  fp32 accumulation (key for NaN prevention on BI-V100)
- Falls back to Python version if .so not available
2026-08-13 06:25:09 +00:00
Claude
c720cbc3a3 docs: dlopen SO开发计划 — gap analysis from comp168 log + real tree 2026-08-13 06:23:59 +00:00
project6-dev
17fdf7e2d6 diag: probe KV cache layout with 5D key + 4D value 2026-08-13 05:22:05 +00:00
project6-dev
cb03fc9993 diag: cat ixformer vllm.py source 2026-08-13 05:20:01 +00:00
project6-dev
b0ed88e114 diag: probe ixformer KV cache 5D layout + read vllm.py source 2026-08-13 05:18:50 +00:00
project6-dev
a8f0332e1c diag: verify_paged_attn.py — test ixformer paged attention v1/v2 with head_dim=256
Now have correct signature: needs head_mapping tensor for GQA.
Tests v1 (basic decode), v2 (partitioned for long context), and performance.
2026-08-13 05:16:25 +00:00
project6-dev
ce568f94ed diag: probe ixformer paged attention signature 2026-08-13 05:14:31 +00:00
project6-dev
ad6863ed84 perf: replace Python Q-tiling fallback with ixformer.flash_attn_varlen_func
Verified on real BI-V100:
  flash_attn_func works with head_dim=256 (diff < 0.004, no NaN)
  flash_attn_varlen_func works for variable-length batching
  seq=1024: 1.7x faster than PyTorch matmul

The profiling-stage _run_sdpa_fallback now tries flash_attn_varlen_func
first, falls back to Python Q-tiling only on exception.

This addresses the 10-50x attention slowdown identified in the analysis:
  Python Q-tiling: O(L^2) per-tile matmul in Python loop
  flash_attn: fused kernel, O(L) memory, hardware-optimized
2026-08-13 05:14:08 +00:00
project6-dev
9f02200ede diag: verify_flash_attn.py — test flash_attn_func correctness + perf for head_dim=256
flash_attn_func WORKS with head_dim=256 on BI-V100!
This is the path to 10-50x attention speedup.
Tests: correctness vs ref, GQA, long seq, varlen, paged decode, perf.
2026-08-13 04:29:49 +00:00
project6-dev
9c97a24edf diag: verify_ixformer_attn.py — test ixformer native attention with head_dim=256
The 10-50x slowdown is from bypassing ixformer SDPA and using Python
matmul fallback. Test if ixformer actually crashes on head_dim=256
or if the bypass was premature.
2026-08-13 04:25:09 +00:00
project6-dev
1aa2262a2c diag: test_triton.py — check if Triton works on BI-V100 2026-08-13 04:18:15 +00:00