Skip to content

2.0.17

Latest

Choose a tag to compare

@krrishnarraj krrishnarraj released this 07 Jul 16:44
96ca6e4

CUDA — WMMA / MMA performance overhaul

  • Native mma.sync kernels — adds direct PTX mma.sync tiles for:
    • FP16 (m16n8k16) — replaces old WMMA fragments, reaches full-rate on Blackwell+
    • BF16 (m16n8k16) — native tile path for BF16 GEMM
    • TF32 (m16n8k8) — replaces the old WMMA m16n16k8 path, matches cuBLASLt peak (~20.6 TFLOPS on Blackwell)
  • FP16 full-rate kernel (wmma_fp16_f16.cu) — fp16 input × fp16 accumulation for full tensor-core rate
  • Dead-code fold fixes — WMMA accumulator chains no longer get optimized away, fixing throughput inflation across FP16, BF16, TF32, INT8 (k32), and FP64

oneAPI — Intel Arc A770 (DG2) joint_matrix

  • Enabled joint_matrix on Intel Arc A770 (DG2) via dispatchJm mechanism
  • Code review cleanups: dtype matching, fallback removal

Metal — simdgroup_matrix dead-code fix

  • Fixed 4× throughput inflation in simdgroup_matrix kernels (accumulator dead-code folding) for FP16 and BF16

OpenCL — Work-group clamping fix

  • Per-kernel work-group size clamping to device max, resolves #183

Android

  • Fixed JNI leaks

Refactoring & infrastructure

  • TextLogger moved to shared common moduleLoggerText in include/common/logger_text.h / src/common/logger_text.cpp. Same formatted output now drives both the desktop CLI and Android (ostream injection), removed the old src/cli/logger_cli.cpp

What's Changed

  • refactor: move TextLogger from CLI to shared common module by @krrishnarraj in #182
  • fix(oneapi): dynamic joint_matrix tile shapes and sub-group size for Intel XMX GPUs by @savvadesogle in #185
  • feat(oneapi): benchmark all device-reported joint_matrix combinations dynamically by @savvadesogle in #186
  • Merge dev → master by @krrishnarraj in #184

New Contributors

Full Changelog: 2.0.16...2.0.17