問題在哪裡?
別人就都沒問題 怎麼就你問題最多??
你怎麼不找找你自己的問題??

![]

meme source:一起來找找問題吧~

P.S
AI 好厲害 原本還想自己做遊戲 結果現在AI隨便做都比你認真學怎麼做遊戲厲害
哈哈 變成玩法 機制更重要了
完蛋 已經意識到 自己要變成那種 開始厭惡新科技的年紀了
哈哈 阿哈哈
離題了，回到AI slob

前兩天我們都在看 native FWT。

Day 16 看 single-block：資料留在 register、warp shuffle、shared memory swizzle。

Day 17 看 large-N：一個 block 放不下，就用 Kronecker split 做 multi-pass。

今天回到一開始那個誘惑：

Hadamard matrix 只有 +1 / -1，
那是不是很適合 Tensor Core？

直覺上是。

實作上很麻煩。

而且更重要的是：就算實作出來，也不一定比較快。
(這真的蠻哭的 就算把形狀改好 還是很靠北)

Tensor Core 想吃的是矩陣乘法

Tensor Core 的 MMA 指令喜歡固定形狀的 tile。

例如 INT8 path 會看到這種指令：

mma.sync.aligned.m16n8k32.row.col.s32.s8.s8.s32

意思是：拿 INT8 的 A tile 和 B tile 做矩陣乘法，累積到 INT32 accumulator。

這對 GEMM 很自然。

但 FWT 的原始形狀不是 GEMM。

FWT 是 butterfly network：

stage 0: local pair add/sub
stage 1: stride 2 add/sub
stage 2: stride 4 add/sub
...

所以要用 Tensor Core，就必須先把一段 Hadamard transform 改寫成 block GEMM。

這就是成本開始出現的地方。

Kronecker block GEMM 的想法

Hadamard 有 Kronecker 結構：

H_N = H_a ⊗ H_b

所以可以把 H_N x 拆成幾個較小的 dense transform。

例如把一段資料看成矩陣，先做一個方向：

Z = X * H_b

再轉置或 reshape，做另一個方向：

Y = H_a * Z

如果 H_a 或 H_b 是小矩陣，元素又只有 +1/-1，就可以想辦法把它塞進 Tensor Core。

case_3 裡的 Tensor Core FP32/TF32 path 大致是這個方向：

reshape input
做 block Hadamard GEMM
transpose intermediate
再做 block Hadamard GEMM
transpose back

聽起來合理。

但 FWT 的問題是，每一步額外搬資料都很貴。

TF32 path：少了 CRT，但還是有 layout 成本

先看比較單純的 TF32 / FP32 Tensor Core path。

這條路沒有 CRT reconstruction，也沒有多 modulus residue。

它主要要處理的是：

1. 把 input 變成 Tensor Core 喜歡的 tile layout
2. 準備 Hadamard block
3. 做 MMA
4. 把中間結果 transpose
5. 做第二次 MMA
6. transpose 回 output layout

在 src/ImplicitHadamardTensorCoreFP32.cu 裡，可以看到 fused path 想把 transpose 放進 kernel 裡處理：

execute_implicit_hadamard_tensor_core_fp32(..., transpose_batch=true)

這個方向是對的，因為 Python 層或 global memory 上額外做 transpose 會很傷。

但即使 fuse 掉一部分，Tensor Core path 還是必須面對一個事實：

FWT 原本只需要加減和資料交換，
現在為了用 MMA，多了 layout conversion。

對 GEMM 來說，layout conversion 可以被大量乘加攤掉。

對 FWT 來說，不一定。

INT8 / Ozaki path：Hadamard 很適合，input 不適合

Hadamard block 本身很適合 INT8：

+1 -> int8
-1 -> int8

問題在 input state。

如果 input 是 FP32 或 FP64，它不能直接丟進 INT8 MMA。你要先轉成整數表示。

前面 GEMM 系列我們做過類似事情：

floating value
-> scaled integer
-> residue mod p
-> balanced int8
-> INT8 GEMM
-> CRT reconstruction

套到 Hadamard 後，Hadamard matrix 的那一邊很簡單，因為它只有 +1/-1。

但 input 這邊仍然要做：

scale
mod p
pack int8
多個 modulus 重複做 MMA
reconstruct

這就很尷尬。

FWT 本身的加減太便宜，Ozaki conversion 的固定成本反而很大。

實作上的資料流

Tensor Core / Ozaki path 大致會變成：

global FP input
-> scale / residue conversion
-> A_mod[p]
-> H_mod[p] 或 implicit Hadamard sign
-> mma.sync int8
-> int32 accumulator
-> mod reduction
-> CRT/scaling
-> output

如果是一般 dense GEMM，A_mod[p] * B_mod[p] 的 K 很大，MMA 會做很多有效工作。

但 Hadamard FWT 不一樣。

Hadamard 的結構本來可以不用 materialize H，也不用做很多無意義的 dense multiply。

所以 Tensor Core path 要成功，必須滿足一個很嚴格的條件：

MMA 省下的時間 > conversion + layout + transpose + reconstruction 的時間

在 case_3 這個問題上，實測不容易成立。

為什麼理論 TFLOPS 會騙人

Tensor Core 的峰值吞吐量很漂亮。

但 FWT 的 arithmetic intensity 太低。

一個 butterfly 只有：

a + b
a - b

對每個 element，每個 stage 只做一次非常便宜的運算。

所以如果你多搬了一次 global memory，成本可能已經比那些加減法還大。

這就是為什麼這題不能只看：

Tensor Core 有多少 TFLOPS

而要看：

用了 Tensor Core 之後，多了多少 bytes traffic？

尤其是 transpose。

一旦 access pattern 不連續，GPU 可能為了拿少量有效資料，搬很多 memory sector。這會把 Tensor Core 的理論優勢直接吃掉。

case_3 的幾個實驗判斷

報告裡有三個很重要的觀察。

第一，傳統 dense Ozaki 對這題太重：

Native FP64 約 600 us
Dense Ozaki GEMM 約 314,000 us

這不是小幅慢，而是方向錯了。

第二，Radix-128 / Kronecker Ozaki 可以把成本拉回來：

Radix-128 Ozaki 約 900 us
Native FP64 約 600 us

這代表 Kronecker 結構有用，但還不夠贏過 native butterfly。

第三，TF32 Tensor Core 少了 CRT 成本，但仍然被 transpose/layout 影響：

Native FP32 約 40 us
TF32 Tensor Core + transpose 約 380 us

這個差距說明：問題不是只有 Ozaki conversion。

就算沒有 CRT，為了餵 Tensor Core 而做的資料重排也可能太貴。

什麼時候 Tensor Core 可能回來

這裡不是說 Tensor Core 對量子模擬沒用。

而是 FWT 這個 operator 太特殊。

如果換成更 dense 的局部 gate、tensor contraction、或比較大的 block operator，情況會不同。

Tensor Core 比較適合：

資料載入後可以重複使用很多次
每個 output 有大量 accumulation
layout conversion 可以被大量 MMA 攤掉

FWT 則剛好相反：

每層只有加減
資料交換比運算重要
不該 materialize Hadamard matrix

所以這篇的結論不是「不要用 Tensor Core」，而是：

不要為了用 Tensor Core，把低 arithmetic intensity 的問題改寫成更重的問題。

Day 18 的結論

Hadamard 的 +1/-1 讓 Tensor Core path 看起來很誘人。

但真正難的不是把 +1/-1 塞進 INT8。

難的是 input layout、transpose、conversion、reconstruction，以及這些成本能不能被 MMA 吞吐量攤掉。

在 FWT 這個問題上，native butterfly 太便宜、資料交換太關鍵，所以 Tensor Core 不一定有主場。

下一篇我們把比較標的和 profile 拉出來看：到底要怎麼公平比較，NCU 又告訴我們什麼。

主委加碼:

寫個程式 讀個資工
寫一個月 阿 AI又更新了
別人怎麼都百萬年薪　隨便AI指令一下就有了
要你幹嘛?
沒有工作怎麼不找找自己問題?
阿叫AI做就好了啊 那麼厲害 都給AI做就好啦
找自己問題
你怎麼不找找你的問題
你怎麼不找找山姆奧特曼的問題?
是我發明AI的嗎?
是我讓你失業的嗎?
你告訴我 怎麼打?(代碼)
GPT又發布新版本拉
然後黃仁薰就在那邊
阿哈哈哈哈 未來不需要工程師 大家應該去當水電工
然後中國又出個什麼迪譜吸客 v4pro v4 破麻 v5 v6.7 v8.7
阿啊啊啊 v8964
然後瑪斯克又在那邊 阿未來人類不用工作
nm了個臭b的 你告訴我怎麼贏? (影片原話)
阿?
你排一群靠AI販賣焦慮衝股票的給我
怎麼贏?
被玩到死 內捲到死
然後沒工作了
親戚就來一句
哎呀 你怎麼不多多找找自己問題呢?
ni ma der
meme source:原影片~

(純粹搞搞抽象 請勿當真 也請尊重每位媽媽跟女性 peace and love)
已閱讀 5 4 6 3 字