Commit Graph

  • 007c645f87 fix(megamoe): normalize fp4 weight preparation contract cch/deepgemm-glm52-nvfp4-group16 LuminolT 2026-07-08 18:48:04 +08:00
  • 2c7543130b feat(megamoe): expose fp4 weight preparation helper LuminolT 2026-07-08 18:36:40 +08:00
  • 8ad348fb11 feat(megamoe): add fp4 group16 to group32 requant path LuminolT 2026-07-08 18:34:33 +08:00
  • 79fcfd6abf feat(megamoe): add nvfp4 group16 capability gate LuminolT 2026-07-08 18:29:09 +08:00
  • 0ff91f1285 docs: record sm90 megamoe phase6 dev log Xinyi Liu 2026-06-18 17:42:09 +08:00
  • 453fc7b046 feat: implement sm90 megamoe phase6 combine Xinyi Liu 2026-06-18 17:40:49 +08:00
  • a0b3bb0017 docs: record sm90 megamoe indentation fix log Xinyi Liu 2026-06-18 15:54:59 +08:00
  • 71f012a356 style: fix sm90 megamoe dispatch indentation Xinyi Liu 2026-06-18 15:54:00 +08:00
  • 6b33b41c4a docs: record sm90 megamoe phase5 dev log Xinyi Liu 2026-06-18 15:18:28 +08:00
  • 9bd0519605 feat: implement sm90 megamoe phase5 l2 scatter Xinyi Liu 2026-06-18 15:17:20 +08:00
  • fc8218750c docs: record sm90 megamoe phase4 dev log Xinyi Liu 2026-06-18 01:11:09 +08:00
  • f3553f976c feat: implement sm90 megamoe phase4 l1 epilogue Xinyi Liu 2026-06-18 01:09:45 +08:00
  • 2bb1756787 docs: record sm90 megamoe phase3 dev log Xinyi Liu 2026-06-18 00:48:12 +08:00
  • 842124b1c5 feat: implement sm90 megamoe phase3 l1 wgmma Xinyi Liu 2026-06-18 00:46:41 +08:00
  • ed6067c7a5 docs: record sm90 megamoe phase2 dev log Xinyi Liu 2026-06-18 00:01:30 +08:00
  • 540e5aeadc feat: implement sm90 megamoe phase2 dispatch-only Xinyi Liu 2026-06-18 00:00:36 +08:00
  • 74eb59dfaa docs: record sm90 megamoe phase1 dev log Xinyi Liu 2026-06-17 23:56:31 +08:00
  • 1dafd0f173 feat: add sm90 megamoe phase1 interfaces Xinyi Liu 2026-06-17 23:54:49 +08:00
  • 062cb160cf Phase 0: SM90 MegaMoE design doc, reference baseline, nsys script Xinyi Liu 2026-06-16 18:01:12 +08:00
  • 714dd1a4a9 Update test_mega_moe.py Chenggang Zhao 2026-05-11 19:20:18 +08:00
  • 891d57b4db Add various optimizations and Mega MoE benchmarks (#316) Zhean Xu 2026-04-24 18:41:37 +08:00
  • 7f2a703ed5 [Public release 26/04] Introducing Mega MoE, FP4 Indexer and other features/fixes (#304) Chenggang Zhao 2026-04-17 09:45:14 +08:00
  • d30fc36c8f Fix sync issue of TMEM alloc/dealloc (#292) Ray Wang 2026-03-22 16:41:28 +08:00
  • 35c4bc8771 fix: k_grouped_fp8_gemm_nt_contiguous crashes with n = 768 on H100 (#238) Xin Qiu 2026-02-25 10:13:54 +08:00
  • 477618cd51 Fix a sync issue in SM100 MQA logits (#285) Ray Wang 2026-02-03 17:29:49 +08:00
  • 0f5f266202 Multiple updates and refactorings (#280) Zhean Xu 2026-01-16 17:06:52 +08:00
  • 3ccf40c53a Merge pull request #270 from yurekami/fix/sm90-archspec-bug Zhean Xu 2026-01-06 09:56:33 +08:00
  • 6be0eb31d9 fix: use SM90ArchSpec instead of SM100ArchSpec in sm90_bf16_k_grouped_gemm yurekami 2026-01-01 05:06:36 +09:00
  • 9b680f4284 Update install.sh Chenggang Zhao 2025-12-05 17:06:48 +08:00
  • 659a279bbd Better error handling, code consistency, compile-time safety (#234) AJ WISE 2025-12-05 08:49:52 +00:00
  • 38f8ef73a4 Multiple updates and refactorings (#231) Ray Wang 2025-11-21 17:49:47 +08:00
  • bb4424aad4 Fix sum_k * shape_m overflow Zhean Xu 2025-11-19 11:51:36 +08:00
  • 8da33d6bd9 Clean up Chenggang Zhao 2025-11-19 11:00:55 +08:00
  • f63d7f24d6 fix: prevent int32 overflow in k-grouped GEMM size calculations (#226) Guoteng 2025-11-19 10:52:08 +08:00
  • ec5e9ed0b8 Fix SM90 MQA logits (#229) Ray Wang 2025-11-19 10:50:36 +08:00
  • 2f9d87877e Use larger MMA shape (#227) Ray Wang 2025-11-14 11:38:15 +08:00
  • c9f8b34dcd Merge pull request #220 from ko3n1g/ko3n1g/chore/revert-name-change oliver könig 2025-10-15 16:30:23 +02:00
  • 237041a257 revert oliver könig 2025-10-15 14:29:57 +00:00
  • f82018273d chore: Revert name change oliver könig 2025-10-15 14:29:16 +00:00
  • 737e420fad chore: Rename project to ds-deem-gemm oliver könig 2025-10-15 12:44:21 +00:00
  • 2b8a8e24f8 Update publish.yml oliver könig 2025-10-15 13:00:51 +02:00
  • 9528451969 Ko3n1g/chore/rename to deepgemm (#217) oliver könig 2025-10-15 12:13:42 +02:00
  • 93b3c28fa8 ci: Fixes for pre-built wheels (#214) oliver könig 2025-10-14 07:05:47 +02:00
  • f8f41145da Use CUDA runtime API to get device prop instead of ATen Chenggang Zhao 2025-10-11 09:14:00 +08:00
  • 9f196058ae chore: Build and store bdist wheels (#181) oliver könig 2025-10-10 12:23:40 +02:00
  • 6e74faad5c Upgrade to CUTLASS 4.2.1 (#203) Jun Jiang 2025-10-09 09:09:22 +08:00
  • 239112cb4c Fix syntax errors and correct the conditional statements (#206) PGFLMG 2025-10-01 20:31:43 +08:00
  • c1bf4cae4b Fix version Chenggang Zhao 2025-10-01 20:31:27 +08:00
  • 07b82fb8cd Fix old CUDA compatibility Chenggang Zhao 2025-10-01 20:29:15 +08:00
  • 594953acce Update version number Chenggang Zhao 2025-09-29 17:12:21 +08:00
  • 0ed3b949d0 Update README Chenggang Zhao 2025-09-29 17:10:12 +08:00
  • 59f2c07cf2 Add SM100 kernels (#201) Simon Mo 2025-09-29 02:07:28 -07:00
  • 80ceeb2c76 Add SM90 kernels (#200) Chenggang Zhao 2025-09-29 17:00:23 +08:00
  • 904b721731 Update README Chenggang Zhao 2025-09-25 16:27:57 +08:00
  • 3f71de7aa9 Make various updates and fixes (#198) Ray Wang 2025-09-25 16:19:07 +08:00
  • 79f48ee15a Fix multicast bug and optimize masked GEMM (#193) yukuai26 2025-09-12 17:12:27 +08:00
  • ea9c5d9270 Use driver API Chenggang Zhao 2025-08-19 11:29:15 +08:00
  • 51d1e9cdd3 Support compilation with CUDA 13.0 (#174) Rain Jiang 2025-08-26 18:30:08 -07:00
  • 0e49c3353b Refactor compiler version checks and arch flags Chenggang Zhao 2025-08-27 09:26:02 +08:00
  • 3a93f4eb28 Fix B200 cu128 NVCC compilation failed (#173) PGFLMG 2025-08-27 09:07:18 +08:00
  • 9c3783beb2 Fix CUBIN symbol name compatibility Chenggang Zhao 2025-08-26 17:42:11 +08:00
  • 89b4089d24 Update test files in README documentation (#169) ZiTian Zhao 2025-08-25 09:43:10 +08:00
  • 2da871e304 Fix grouped gemms performance issue. (#168) zhonghui-J 2025-08-22 17:35:43 +08:00
  • e38c2e3103 Remove comments Chenggang Zhao 2025-08-22 17:32:04 +08:00
  • f20256fd50 Compatible with CUDA 13 Chenggang Zhao 2025-08-22 17:29:10 +08:00
  • affdb1cd90 Add sm_100f support and make nvcc 13 happy (#157) xiweny 2025-08-22 17:19:32 +08:00
  • f85ec649d7 Make various updates and fixes: (#164) Ray Wang 2025-08-15 18:32:35 +08:00
  • 3254b758e2 Polish get_best_configs modeling. (#158) zhonghui-J 2025-08-14 16:50:21 +08:00
  • 6d3717d541 Update test_fp8.py (#159) fzyzcjy 2025-08-14 16:47:57 +08:00
  • 7b6b5563b9 Fix smxx layout assertion (#154) LJC00118 2025-08-05 10:38:06 +08:00
  • 3979c0576e Merge pull request #151 from RayWang96/update_jit Ray Wang 2025-08-03 11:04:02 +08:00
  • d9c363f86f Make various updates and fixes: - Add support for legacy CUDA versions; now compatible with CUDA 12.3 and newer - Add support for NVRTC compilation - Other fixes and code refactoring Ray Wang 2025-08-02 19:52:22 -07:00
  • aff9da0aba Fix SM90 GEMM (#149) yukuai26 2025-08-01 10:36:49 +08:00
  • c50deed14c Code lint Chenggang Zhao 2025-07-30 10:39:30 +08:00
  • 6bc75b549e Fix smxx layout assertion (#141) LJC00118 2025-07-30 10:36:54 +08:00
  • a581263f06 Fix indent LJC00118 2025-07-29 16:14:46 +08:00
  • fb7c687548 Merge pull request #135 from danthe3rd/patch-3 dan_the_3rd 2025-07-29 10:14:15 +02:00
  • dd6ed14acb Add torch as build dependency. (#139) Yuxian Qiu 2025-07-28 17:01:39 +08:00
  • 4b4e4f20dd Update system.hpp (#133) dan_the_3rd 2025-07-28 11:01:05 +02:00
  • 8987798502 Update setup.py (#134) dan_the_3rd 2025-07-28 10:58:23 +02:00
  • 187656694f Code lint Chenggang Zhao 2025-07-21 11:00:50 +08:00
  • 436a56314c Use std::filesystem::directory_iterator instead of std::filesystem::recursive_directory_iterator to avoid an ABI breakage we met (#131) Ray Wang 2025-07-21 10:44:20 +08:00
  • c1db17e2e3 Updated submodules to use https:// vs git@ (#129) Clayton Coleman 2025-07-20 21:05:15 -04:00
  • 4ca3cdffd2 fix: update .gitmodules (#130) Yineng Zhang 2025-07-19 19:23:58 -07:00
  • 6c9558ed04 Update CUDA toolkits requirement (#128) Ray Wang 2025-07-18 16:40:47 +08:00
  • 9da4a23561 Add more GPU architectures support (#112) Ray Wang 2025-07-18 11:32:22 +08:00
  • 03d0be3d2d Simplify expression Chenggang Zhao 2025-07-02 14:07:05 +08:00
  • 3fc6728dee [add] fix smem_barrier size in wgrad way (#122) fy1214 2025-07-02 14:05:36 +08:00
  • e82c4139da Revert "Fixed the bug in get_swizzle_mode function related to elem_size setting. (#115)" yukuai 2025-06-23 17:13:36 +08:00
  • ac428e25e0 Fixed the bug in get_swizzle_mode function related to elem_size setting. (#115) TherLF 2025-06-23 09:37:10 +08:00
  • 0c88cd0139 Fix illegal memory address when skipping -1 m indices (#113) shixianc 2025-06-15 19:44:31 -07:00
  • 8dfa329827 Grouped GEMM skip useless computation for unaligned Ms (#103) yukuai26 2025-05-27 13:43:38 +08:00
  • 391755ada0 Fix JIT tests Chenggang Zhao 2025-05-16 14:39:58 +08:00
  • 78d8362e7a Add a missing #pragma once Chenggang Zhao 2025-05-15 18:10:05 +08:00
  • ec426b9d66 Merge pull request #100 from deepseek-ai/remove-tuner Chenggang Zhao 2025-05-15 17:05:42 +08:00
  • 104a6ec109 Add __assertfail Chenggang Zhao 2025-05-15 17:04:21 +08:00
  • 3b412f458a Unify kwargs usages Chenggang Zhao 2025-05-15 16:53:52 +08:00
  • 350989eef3 Unify ceil_divs Chenggang Zhao 2025-05-15 16:48:32 +08:00
  • 4373af2e82 Add DG_PRINT_CONFIGS Chenggang Zhao 2025-05-15 16:36:40 +08:00
  • 816b39053a Refactor launch-related structures Chenggang Zhao 2025-05-15 16:14:21 +08:00