Skip to content

feat: update grouped GEMM, fused MoE, and convolution support - #6

Merged
longbisheng merged 11 commits into
mainfrom
sync_dcudas
Oct 19, 2026
Merged

longbisheng merged 11 commits into
mainfrom
sync_dcudas

Conversation

@longbisheng

Copy link
Copy Markdown
Contributor

Adds grouped GEMM tile-loop, AB-scale and WASP pipelines, integrates BF16 NN/TN kernels as CK instances, and adds fused MoE and FP32 convolution support. Includes GEMM argument checks and compiler compatibility fixes. Preserves existing public CI files and all 11 individual commits.

Validation: per-commit metadata and tree verification, complete Linux checkout, and strict Git object checks passed. GPU compilation, correctness, and performance tests were not rerun for this synchronization.

Please merge while preserving the branch commit IDs. Dependent AITER snapshots use exact public commit mappings; squash or rebase would require reconciling those mappings.

wangwanqi and others added 11 commits September 29, 2026 20:25
… CK instances

Port the standalone uniform Direct-to-LDS NN (dgrad) / TN (wgrad) kernels
into the ck_tile framework as device-args instances:

- gemm_pipeline_v7_uniform.hpp: uniform (non warp-specialized) pipeline,
  256x256x16, 512 threads, 4-stage 16KB Direct-to-LDS ring, two-group wave
  phase interleave; V7UniformKind::NN walks a K-contiguous A, ::TN walks
  the [K, M] panel (both operands m32x16).
- grouped_gemm_v7_uniform_epilogue.hpp: per-wave LDS slab staging, 16B
  vector stores, no cross-warp barrier.
- instances 5000 (v7_bf16_nn) / 5001 (v8_bf16_tn), family v7, opt-in via
  CK_TILE_GROUPED_GEMM_ENABLE_BF16_V7_NN / _V8_TN (disable wins); selector
  gates: gfx936, bf16, K>=64 && K%16==0, N%256==0, G<=16, fixed K, no
  empty groups.
- entries are NOT registry candidates: a dedicated archive carries the
  tuned codegen contract (-mllvm -support-768-vgprs et al.); the flat
  over-covered grid early-exits blocks past the last group's block_end
  (FindGroupId only converges in-range), and the LDS ring is DYNAMIC
  shared memory addressed from offset 0 per the reference kernels.
- shared-header decoupling: kernel.hpp only includes the 10-line forward
  header (complete tile-parameter POD + enum); the pipeline body and the
  GemmConfigV7Uniform configs compile only inside the v7/v8 TUs, so
  tuning recompiles two translation units, not the archive.
- grouped_gemm_v7_instance_test: correctness (CPU reference + default
  instance triangle) and perf through the versioned device-args ABI;
  selector/query tests extended to the 20-candidate registry.

gfx936, groups=4 M=[8100,8100,8091,8091] N=4096 K=4096:
  NN 2.741 ms / 396.7 TF (default instance 6.508 ms / 167.1 TF)
  TN 2.700 ms / 402.7 TF
Add the CK Wasp producer-consumer pipeline, grouped single-group launch path, direct-store epilogue, and focused correctness/performance driver. Preserve the C/Cp stage schedule while hoisting producer wave-ID reads out of the K loop.

Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
@hygon-ai-reviewer

hygon-ai-reviewer Bot commented Oct 19, 2026 •

Copy link
Copy Markdown

AI Review

发现 4 个问题。

[P2] conv 示例 pad 参数向 ConvParam 传入错位

  • 位置:example_hcu/ck_tile/14_conv/conv_fwd_fp32.cpp(行号未确定)
  • 触发条件:用户通过命令行设置不相等的 pad h 与 pad w(例如 p=2、q=1)运行该示例。
  • 证据与影响:ArgParser 中 p 定义为 pad h、q 定义为 pad w,但构造 conv_param 时 input_left_pads_ 传 {p, p}、input_right_pads_ 传 {q, q};按参数语义应为两处均传 {p, q}。 p 与 q 不相等时,padding 会按错误的维度和错误的对齐侧应用,参考卷积与设备卷积按错误输入形状计算,校验结果失真或示例异常失败。
  • 修改建议:将 input_left_pads_ 与 input_right_pads_ 两处均改为 {arg_parser.get_int("p"), arg_parser.get_int("q")},并人工核验 ConvParam 的维度顺序约定。

[P3] 标注待删除的 TEMP-DEBUG 探针代码遗留在示例中

  • 位置:example_hcu/ck_tile/17_fused_moe/fused_moe_flatmm_uk_fp16.cpp(行号未确定)
  • 触发条件:每次运行该示例时,内核启动前必然执行该调试代码块
  • 证据与影响:代码注释明确写有 "TEMP-DEBUG probe (remove after diagnosis)",随后是多条 printf("DBG: ...") 调试输出,以及对 d_a/d_g/d_d 各做一次 16 字节 device-to-host 回读 示例每次运行都向 stdout 固定输出调试信息,并引入不必要的设备往返拷贝;作为随仓库发布的参考示例会误导使用者
  • 修改建议:合并前删除整个 TEMP-DEBUG 调试块(printf 输出与回读拷贝),仅保留内核参数设置与启动逻辑

[P3] 计时区间内包含 hipMemset,测得的性能数据包括清零开销

  • 位置:example_hcu/ck_tile/17_fused_moe/fused_moe_flatmm_uk_fp16.cpp(行号未确定)
  • 触发条件:任何以 repeat>0 运行该示例的性能测量路径
  • 证据与影响:hipEventRecord(e0) 之后、hipEventRecord(e1) 之前的循环体内,每次迭代先执行 hipMemset(d_o, 0, o.size() * 4) 再 launch(),随后对 e0..e1 求得的耗时除以 repeat 打印为 time/TFlops 报告的 time 与 TFlops 系统性地包含每次迭代的设备端清零耗时,输出性能数值偏离 kernel 实际耗时,影响基准对比参考价值
  • 修改建议:将 memset 移出计时区间(如仅在计时前清零一次,或以两次事件仅包围 launch 调用),确保计时段只含 kernel 启动与执行

[P3] 退化输出早退分支跳过所有 hipFree,泄漏设备显存

  • 位置:example_hcu/ck_tile/17_fused_moe/fused_moe_flatmm_uk_fp16.cpp(行号未确定)
  • 触发条件:参考或输出存在非有限值或退化(nonfinite_o>0、nonzero_ref==0 等成立)时进入 FAIL 分支执行 return ret。
  • 证据与影响:该分支内 ret=1 后直接 return ret,而 d_a、d_g、d_d、d_o、d_sid、d_sw、d_seid、d_nst 的 hipFree 全部位于函数末尾正常路径,早退路径不会执行。 错误路径上已分配的设备缓冲泄漏至进程结束;若示例在单进程中多次运行不同配置且命中失败,显存占用累积,可能影响后续分配。
  • 修改建议:将 hipFree 提前到 return 前统一清理,或改为单一出口/RAII 方式,确保所有错误分支都释放设备缓冲。

审查未完整完成:以 GitHub 返回的 PR diff 为主体,关联上下文仅用于核验,未运行代码。

覆盖情况:共 66 个文件,已审查 11 个,待审查 54 个,跳过 1 个。

跳过原因:无 patch 1 个;文件类型过滤 0 个。

进度已保存。正在继续处理。

@longbisheng
longbisheng merged commit e295c92 into main Oct 19, 2026
4 checks passed
@longbisheng
longbisheng deleted the sync_dcudas branch October 19, 2026 07:12
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants