add bf16 exp fallback (#776)

xwhzz · web-flow · commit 471cc7f846fe · 2025-09-01T20:59:55.000+08:00
diff --git a/src/tl_templates/cuda/common.h b/src/tl_templates/cuda/common.h
@@ -330,3 +330,8 @@ TL_DEVICE void __sync_thread_partial() {
   asm volatile("bar.sync %0, %1;" : : "r"(barrier_id), "r"(thread_count));
 }
 } // namespace tl
+
+namespace cutlass {
+TL_DEVICE
+bfloat16_t fast_exp(bfloat16_t x) { return ::hexp(x); }
+} // namespace cutlass

Original file line number	Diff line number	Diff line change
`@@ -330,3 +330,8 @@ TL_DEVICE void __sync_thread_partial() {`
`330`	`330`	`asm volatile("bar.sync %0, %1;" : : "r"(barrier_id), "r"(thread_count));`
`331`	`331`	`}`
`332`	`332`	`} // namespace tl`
	`333`	`+`
	`334`	`+namespace cutlass {`
	`335`	`+TL_DEVICE`
	`336`	`+bfloat16_t fast_exp(bfloat16_t x) { return ::hexp(x); }`
	`337`	`+} // namespace cutlass`