
KurisandLeiWang1999
719d23dce5
feat: auto-vectorize bf16/fp16 reduce with packed add2 intrinsics (#2112)
* fix: add missing TvmLogDebugSettings::ParseSpec and VerboseEnabledImpl for TVM_LOG_CUSTOMIZE builds
When TVM_LOG_CUSTOMIZE=1, TVM's logging.cc skips compiling
TvmLogDebugSettings::ParseSpec and VerboseEnabledImpl (guarded by
#if TVM_LOG_CUSTOMIZE == 0). However libtilelang.so calls these
functions via the inline FromFlag(), causing a runtime symbol
lookup error. Add the missing implementations to tilelang's own
logging.cc.
* feat: auto-vectorize bf16/fp16 reduce with packed add2 intrinsics
* style: clang-format
* fix: fp16 batch pack — fix Shuffle codegen taking address of temporary
- Add tl::pack_half2 helper in common.h that returns uint1, avoiding taking
address of a temporary __pack_half2 return value
- Update ShuffleNode codegen for fp16x2 to use pack_half2 instead of
&(__pack_half2(...))
- Re-enable fp16 batch pack (remove bf16-only restriction in can_batch_pack)
- Add SumOp_fp16x2/MaxOp_fp16x2/MinOp_fp16x2 packed reducer structs
* refactor: simplify reduce packing API
- Replace ReducePackConfig with simple GetPreferedVectorizedSize(DataType, Target) -> int
- Merge getBatchReducerSuffix into MakeCodegenReducer(vsize)
- MakeReduce(vsize, acc, b) -> optional<PrimExpr>
- MakeInitValue(vsize) handles Broadcast internally
- Guard nan_propagate in MakeCodegenReducer vsize>1 path (no packed nan ops exist)
* fix: guard nan_propagate in MakeReduce packed path, mirror scalar structure
* feat: support nan_propagate in packed reduce (max2_nan/min2_nan)
* fix: add bitwise reduce ops to MakeReduce vsize=1 path
* fix lint error
---------
Co-authored-by: LeiWang1999 <leiwang1999@outlook.com>
2026-05-19 15:03:44 +08:00
..
2025-12-12 17:22:18 +08:00
2026-04-20 16:35:43 +08:00
2026-02-11 22:51:59 +08:00
2026-01-06 02:18:29 +08:00
2026-04-07 23:54:55 +08:00
2026-03-27 11:17:21 +08:00
2026-03-23 19:26:24 +08:00
2026-01-06 02:18:29 +08:00
2026-04-07 23:54:55 +08:00
2026-04-07 23:54:55 +08:00
2025-12-17 14:41:25 +08:00
2026-03-06 17:35:16 +08:00
2026-05-09 16:04:31 +08:00
2026-05-12 15:13:34 +08:00
2026-02-09 17:30:22 +08:00
2025-12-17 14:41:25 +08:00
2026-04-07 23:54:55 +08:00
2025-12-17 14:41:25 +08:00
2026-04-21 18:27:15 +08:00
2026-04-17 15:19:04 +08:00
2026-04-07 23:54:55 +08:00
2026-01-04 15:45:38 +08:00
2026-04-07 23:54:55 +08:00
2026-03-18 13:08:30 +08:00
2026-05-10 21:50:30 +08:00
2026-03-18 13:08:30 +08:00
2026-03-18 13:08:30 +08:00
2026-04-02 12:42:47 +08:00
2025-12-17 14:41:25 +08:00
2026-04-13 12:42:44 +08:00
2026-01-06 02:18:29 +08:00
2026-05-06 23:08:24 +08:00
2025-12-17 14:41:25 +08:00
2026-03-12 17:21:01 +08:00
2026-02-02 18:47:18 +08:00
2026-04-07 23:54:55 +08:00
2026-04-21 18:09:24 +08:00
2026-04-07 23:54:55 +08:00
2026-04-13 12:42:44 +08:00
2025-12-17 14:41:25 +08:00
2026-03-18 13:08:30 +08:00
2026-01-09 17:31:33 +08:00
2026-04-07 23:54:55 +08:00
2026-04-21 18:27:15 +08:00
2026-01-29 23:59:24 +08:00
2026-04-13 16:18:24 +08:00
2026-05-19 15:03:44 +08:00
2026-05-08 17:11:34 +08:00
2026-02-02 18:47:18 +08:00
2026-04-17 16:48:28 +08:00
2026-02-02 18:47:18 +08:00
2026-04-14 15:54:28 +08:00
2026-05-18 14:50:43 +08:00
2026-05-06 23:08:24 +08:00
2025-12-17 14:41:25 +08:00
2025-12-17 14:41:25 +08:00
2026-04-28 12:35:16 +08:00
2026-05-03 23:57:20 +08:00
2026-05-06 23:08:24 +08:00
2026-01-12 22:33:17 +08:00
2026-01-06 02:18:29 +08:00
2026-05-08 17:11:34 +08:00
2026-01-27 20:54:40 +08:00
2026-04-17 17:22:54 +08:00
2025-12-17 14:41:25 +08:00
2026-04-14 01:05:54 +08:00
2026-04-14 01:05:54 +08:00
2026-04-07 23:54:55 +08:00
2026-04-07 23:54:55 +08:00