剪輯：剪輯部

斯坦福和普林斯頓商酌者發(fā)現(xiàn)，DeepSeek-R1生成的自界說CUDA內(nèi)核，完爆了o1和Claude 3.5 Sonnet，拿下總排行第一。誠然面前只可在不到20%任務(wù)上卓絕PyTorch Eager基線，但GPU編程加快自動化的按鈕，照舊被按下！

近日，來自斯坦福和普林斯頓的商酌者發(fā)現(xiàn)，DeepSeek-R1照舊能生成自界說CUDA內(nèi)核了，而且還在一眾推理模子中，徑直拿下了TOP 1！

緊隨后來，OpenAI o1和Claude 3.5 Sonnet鑒識排第二和第三。

具體過程，便是給定一個PyTorch法子，讓模子對其優(yōu)化，然青年景一個包含自界說CUDA內(nèi)核的PyTorch版塊。

在此期間中，模子不錯開脫決定優(yōu)化哪些操作，以提高籌劃效率。

輔導(dǎo)模子生成GPU內(nèi)核，后勁紛亂

如今，傳統(tǒng)東說念主工優(yōu)化內(nèi)核的格式，在效率上照舊不及以應(yīng)酬多數(shù)暴露的AI模子架構(gòu)和硬件平臺。

既然是為了LLM進行優(yōu)化，那么，咱們是否也冒失借助LLM來模擬AI工程師的責(zé)任經(jīng)過，憑借編譯器反應(yīng)、硬件規(guī)格等豐富的信息，自動編寫出準確且經(jīng)過優(yōu)化的內(nèi)核代碼呢？

為此，商酌團隊提議了一種全新的KernelBench框架，用于生成和評估不同AI任務(wù)（單個操作、操作序列、端到端架構(gòu)）的內(nèi)核，并模擬了工程師迭代優(yōu)化的過程。

論文地址：https://arxiv.org/abs/2502.10517

GPU的本體，是硬件依賴的。因此，商酌者們但愿嘗試，看是否能通過以下格式，輔導(dǎo)模子生成GPU內(nèi)核。

最初，向模子提供硬件信息（如內(nèi)存帶寬、TFLOPS），以針對特定GPU（A100、H100等）進行優(yōu)化。

然后，要讓模子在崎嶇文中展示代表性的內(nèi)核優(yōu)化手段，舉例矩陣乘法中的分塊（tiling）或Flash Attention中的在線softmax。

商酌者們發(fā)現(xiàn)，唯一更強勁的模子，會偶爾闡明出欺詐這些優(yōu)化的能力。

崎嶇滑動稽察

比如，DeepSeek-R1未必會使用特定于硬件的指示（如Tensor Core的wmma），但時常無法正確編譯或使用它們，從而脅制了最終性能。

總的來說，商酌發(fā)現(xiàn)，前沿模子在KernelBench上的開箱即用性較差，OpenAI o1和DeepSeek-R1在不到20%的任務(wù)上高出PyTorch Eager基線。

這些模子生成的內(nèi)核存在多數(shù)引申失實、功能正確性問題，何況無法進行特定平臺的優(yōu)化。

具體來說，商酌者發(fā)現(xiàn)：

對模子而言，編寫功能正確的內(nèi)核仍然具有挑戰(zhàn)性；

模子通過優(yōu)化展示了生成高性能內(nèi)核的后勁；

欺詐反應(yīng)關(guān)于減少引申失實和發(fā)現(xiàn)更快的決策很伏擊。

天然，KernelBench面前還只是讓GPU加快奔走的發(fā)軔，但亦然讓通盤GPU編程自動化的肇始催化劑。

令東說念主歡快的是，面前照舊有了好多新的責(zé)任，專注于處理KernelBench中觸及的問題。

比如2月12日，英偉達就發(fā)出博客著述，探討奈何使用DeepSeek-R1進行GPU內(nèi)核自動生成與推理時scaling。

隨后在2月12日，Meta也發(fā)文測試了前沿模子編寫GPU內(nèi)核方面的性能，他們發(fā)現(xiàn)，最好模子不錯在KernelBench上提供平均1.8倍的加快。

Sakana AI更是推出「AI CUDA工程師」，讓AI我方寫代碼優(yōu)化CUDA內(nèi)核，速率宣稱比PyTorch原生完結(jié)快了10-100倍。

如棋布星陳般出現(xiàn)的商酌標明，如今，咱們照舊參加了AI驅(qū)動加快AI的新紀元！

在當年，KernelBench還將捏續(xù)演進。它不會僅限于刻下網(wǎng)羅的250個問題，還不錯彭脹到新的AI任務(wù)。與此同期，評測主張fast_p也不錯跟著時辰的推移進行調(diào)遣，提高加快門檻，以鼓勵更高效的優(yōu)化決策

KernelBench：AI內(nèi)核生成框架

KernelBench是一個開源框架，旨在評估LLM在編寫GPU內(nèi)核方面的能力。

任務(wù)格式

KernelBench包含250個任務(wù)，涵蓋了多樣AI責(zé)任負載，何況易于彭脹到新的責(zé)任負載。

下圖1展示了KernelBench評估講話模子（LM）生成高性能GPU內(nèi)核的能力。KernelBench條件講話模子為給定的標的PyTorch模子架構(gòu)生成優(yōu)化的CUDA內(nèi)核，并進行自動化評估。

· 任務(wù)輸入

給定一個AI責(zé)任負載，任務(wù)的輸入是用PyTorch編寫的參考完結(jié)。效法商酌東說念主員的責(zé)任經(jīng)過，PyTorch代碼包含一個秉承自torch.nn.Module ()的名為Model的類，其中圭臬的__init__和 forward () 函數(shù)（以及任何輔助函數(shù)）被填充為AI責(zé)任負載的PyTorch操作。

AI算法閑居在大型張量數(shù)據(jù)上進行操作。責(zé)任負載的最優(yōu)內(nèi)核取決于張量的大小和數(shù)據(jù)類型（如BF16、FP8）。因此，每個任務(wù)還包含get_inputs ()和get_init_inputs ()函數(shù)，用于指定內(nèi)核需要處理的精準輸入張量。

· 任務(wù)輸出

給定輸入，LLM需要輸出一個秉承自torch.nn.Module ()的名為ModelNew的新類，其中包含自界說優(yōu)化。舉例，LLM不錯在forward ()函數(shù)中使用PyTorch的CUDA-C彭脹來集成內(nèi)聯(lián)內(nèi)核調(diào)用。

為了見效完成任務(wù)，模子需要服氣（1）Model類中的哪些操作最能從優(yōu)化中受益；（2）奈何優(yōu)化這些操作。LLM不錯使用任何硬件高效時代（如交融和分塊）或?qū)Ｓ弥甘荆ㄈ鐝埩恐袠校┮约叭魏尉幊處欤ㄈ鏟TX、CUDA、CUTLASS、Triton、ThunderKittens）。

崎嶇滑動稽察

任務(wù)遴薦

這些任務(wù)左證包含的基本操作或PyTorch庫函數(shù)的數(shù)目分為三個級別。

Level 1包含100個單個基本操作，如卷積、矩陣乘法等AI基礎(chǔ)構(gòu)建塊。誠然PyTorch調(diào)用了經(jīng)過優(yōu)化的閉源內(nèi)核，讓LLM卓絕基線具有挑戰(zhàn)性，但如果能生成開源內(nèi)核，將有伏擊價值。

Level 2包含100個操作序列，如卷積、ReLU和Bias的組合，這些操作不錯交融成一個內(nèi)核以提高性能。

由于基于編譯器的器具（如PyTorch編譯器）在交融方面相稱靈驗，LLM要在這方面卓絕它們也具有挑戰(zhàn)性。然則，LLM可能會提議更復(fù)雜的算法。

Level 3包含50個竣工的機器學(xué)習(xí)架構(gòu)，如AlexNet和MiniGPT等，這些架構(gòu)在運行進修和推理時對內(nèi)核的性能條件極高。

評估主張

KernelBench引入了一個新的評估主張fast_p，估量生成的內(nèi)核中功能正確且加快大于閾值p的任務(wù)比例。

通過調(diào)遣閾值參數(shù)p，商酌者不錯評估不同加快閾值下的內(nèi)核性能，并捕捉加快漫衍。

fast_0終點于LLM的正確率，它估量代碼功能正確的任務(wù)比例，而不研討其速率。在試驗評估中，閑居以p=1動作發(fā)軔。

LLM在KernelBench上的闡明

商酌東說念主員對一系列LLM在KernelBench上進行了評估，收尾自滿，面前的LLM在生成正確且優(yōu)于PyTorch基線速率的內(nèi)核方面仍有艱苦。

在一次性基線評估中，LLM生成的內(nèi)核平均在不到20%的任務(wù)中比PyTorch Eager更快。這標明，僅靠膚淺領(lǐng)導(dǎo)，LLM很難在性能上卓絕傳統(tǒng)的PyTorch內(nèi)核。

LLM生成的內(nèi)核存在多數(shù)的引申失實和功能正確性問題，經(jīng)常由于膚淺的編譯器和運行時失實而失敗。

引申失實包括CUDA/nvcc/Python編譯時失實、CUDA內(nèi)存違法和運行時失實等；正確性失實則主要闡明為輸出張量局勢和值不匹配。

推理模子（o1，R1）生成的失實處理決策（<55%）比其他模子（>70%）少。然則，這主淌若由于引申失敗的情況較少。在功能正確性方面，總計LLM王人瀕臨相似進度的艱苦。

在性能方面，模子生見遵守正確的內(nèi)核在多數(shù)情況下也未能優(yōu)于PyTorch基線。

跟著p的提高，模子生成的內(nèi)核中能達到條件的比例越來越低。在p=1時，在總計KernelBench級別中，不到15%的LLM生成內(nèi)核優(yōu)于PyTorch。

推理模子閑居在提供加快方面優(yōu)于其他LLM，但總體仍有不及。

模子生成的內(nèi)核在不同硬件平臺上的通用性欠安。DeepSeek-R1生成的內(nèi)核在NVIDIA L40S上完結(jié)了36%的加快，而在NVIDIA A10G上則為47%。

這標明LLM在生成特定標的硬件的高效內(nèi)核方面還存在很大的提拔空間。

模子能力分析

測試時欺詐KernelBench環(huán)境反應(yīng)

正如上頭不雅察到的，引申失敗是LM生成的內(nèi)核中最常見的失敗模式。

KernelBench提供的環(huán)境允許網(wǎng)羅豐富的信號，包括編譯器失實、正確性查驗和運行時性能分析主張，總計這些王人不錯反應(yīng)給LM以匡助它處理內(nèi)核故障。

為了探索LM奈何欺詐這些反應(yīng)，商酌團隊評估和比較了兩個基線：第一個令LM為每個KernelBench任務(wù)生成多個并行樣本，另一個通過允許LM欺詐引申反應(yīng)漸漸創(chuàng)新，挨次為每個KernelBench任務(wù)生成內(nèi)核。

· 重迭采樣

KernelBench環(huán)境營救對LM生成的內(nèi)核進行法子化考據(jù)，允許商酌團隊網(wǎng)羅和評估每個任務(wù)的多個LM生成。他們使用fastp@k評估這種重迭采樣門徑。重迭采樣有助于LM發(fā)現(xiàn)更多快速且正確的處理決策。

如下圖4所示，跟著k值的加多，在DeepSeek-V3和Llama 3.1 70B的三個級別上，通過高溫度參數(shù)重迭采樣不錯提拔fast1的性能。

值得珍視的是，在Level 2上，DeepSeek-V3在k=100個樣本時達到了37%的fast1，而在單次領(lǐng)導(dǎo)基線中僅為4%。

通過查驗樣本，咱們發(fā)現(xiàn)高溫采樣有助于探索處理決策空間，加多了生成具有更好優(yōu)化的無失實內(nèi)核的契機。然則，如果一個模子處理任務(wù)的固有概率相稱低，只是加多采樣預(yù)算的影響有限。

舉例，即使嘗試了100個樣本，DeepSeek-V3也從未冒失為Level 1中的一組34個卷積變體生成任何正確的處理決策。

· 生成收尾的迭代優(yōu)化

KernelBench環(huán)境相稱合乎網(wǎng)羅編譯器反應(yīng)、引申失實和使用PyTorch分析器等器具進行的時辰分析，動作果然信號（ground-truth signals）。

商酌東說念主員商酌了欺詐這些反應(yīng)是否能匡助講話模子（LMs）迭代優(yōu)化其生成收尾。

下圖5自滿，KernelBench框架使模子冒失在迭代優(yōu)化過程中領(lǐng)受并欺詐反應(yīng)。這些果然信號包括NVCC編譯器失實信息、引申統(tǒng)計數(shù)據(jù)（舉例正確性查驗和掛鐘時辰），以及PyTorch分析器（操作時辰認識）。

他們在多輪過程中為模子提供每次生成的反應(yīng)：在運轉(zhuǎn)生成后，向模子提供其之前的生成收尾G，以及刻下生成對應(yīng)的編譯器/引申反應(yīng)E和/或分析器輸出P。

然后將每次生成偏激后續(xù)反應(yīng)界說為一輪（turn），并在N輪內(nèi)運行這一迭代優(yōu)化過程。欺詐引申反應(yīng)有助于減少失實，并隨時辰提拔舉座加快效率。

商酌東說念主員不才表2中查驗了第N=10輪時的fast1行為，發(fā)現(xiàn)迭代優(yōu)化在不同模子和KernelBench的各個級別上均捏續(xù)提拔了性能。

DeepSeek-R1在Level 2上的創(chuàng)新最為權(quán)臣，其中引申反應(yīng)E和分析器反應(yīng)P的組合將fast1從36%提拔至72%（如下圖6所示）。

此外，通過分析迭代優(yōu)化軌跡，他們發(fā)現(xiàn)模子在引申反應(yīng)E的匡助下能更靈驗地自我矯正，尤其是在建設(shè)與引申失實籌商的問題上。

DeepSeek-R1在Level 1和Level 2上，經(jīng)過10輪優(yōu)化后，能在高出90%的任務(wù)中生見遵守正確的內(nèi)核（下表9）。

然則，剩余的失實內(nèi)核簡直老是由于功能不正確而失敗，這可能是因為正確性反應(yīng)的顆粒度不如引申失敗信息考究。

· 比較重迭采樣與迭代優(yōu)化

在上表2中，商酌東說念主員比較了在固定10次推理調(diào)用預(yù)算下重迭采樣和迭代優(yōu)化的效率。兩種門徑相較于單次基線（one-shot baseline）均取得了權(quán)臣創(chuàng)新，其中迭代優(yōu)化在6個案例中的5個中闡明更優(yōu)。

然則，他們最終發(fā)現(xiàn)，測試時門徑的效率本體上依賴于基礎(chǔ)模子的質(zhì)地。

舉例，在重迭采樣中，DeepSeek-V3在總計三個級別上恒久優(yōu)于Llama-3.1 70B。相似地，在迭代優(yōu)化中，DeepSeek-R1通過反應(yīng)E和P捏續(xù)創(chuàng)新，而DeepSeek-V3和Llama-3.1 70B并非總能從這些信息中獲益。

提供硬件學(xué)問生成硬件高效內(nèi)核

昭彰，講話模子在生成硬件高效內(nèi)核方面闡明存限。

這可能是由于進修數(shù)據(jù)中內(nèi)核代碼的稀缺性，以及最好內(nèi)核可能需要左證硬件平臺的特定屬性而變化。

在本案例商酌中，商酌團隊探索了提供以下內(nèi)容的效率：（1）提供內(nèi)核工程最好實踐的示例，并將其置于（講話模子的）崎嶇文之中；（2）提供詳備的硬件規(guī)格施展，并將其置于（講話模子的）崎嶇文之中。

· 硬件感知的崎嶇文示例

編寫細密的內(nèi)核閑居使用交融（fusion）、分塊（tiling）、重籌劃（recompute）和異步（asynchrony）等時代來最大化性能。

具體來說，商酌東說念主員納入了三個崎嶇文示例：使用操作交融的GeLU、使用分塊的矩陣乘法，以及展示分享內(nèi)存I/O經(jīng)管的最小Flash-Attention內(nèi)核。

收尾則自滿，崎嶇文示例裁減了講話模子的舉座fast1分數(shù)，因為模子嘗試了更激進的優(yōu)化戰(zhàn)術(shù)，但導(dǎo)致更多引申失敗。與上頭基線生成的代碼比較，OpenAI o1在使用少樣本示例時生成的代碼平均長度加多了25%。

然則，在正確的處理決策中，講話模子應(yīng)用了一些真諦的優(yōu)化：他們發(fā)現(xiàn)，在KernelBench Level 1的77%的GEMM變體中，o1應(yīng)用了分塊并提拔了速率，優(yōu)于單次基線。在Level 2，o1在11個問題上應(yīng)用了激進的分享內(nèi)存I/O經(jīng)管，并冒失卓絕PyTorch Eager。

· 指定硬件信息

正如上頭籌謀過的，內(nèi)核性能因硬件平臺而異。

舉例，F(xiàn)lashAttention-2從NVIDIA A100搬動到H100 GPU時硬件欺詐率下落了47%。FlashAttention-3是一個全王人不同的算法，專為H100編寫。

在本商酌中，商酌團隊探討講話模子是否能欺詐崎嶇文中的以下信息生成創(chuàng)新的內(nèi)核：（1）硬件規(guī)格，舉例GPU類型（H100、A100等）、內(nèi)存大小、帶寬、TFLOPS；（2）硬件學(xué)問（舉例線程、線程束、線程塊、流多處理器的界說）。

收尾自滿，模子很少生成針對底層硬件優(yōu)化的內(nèi)核，這標明當年模子仍有創(chuàng)新空間。

某些新一代GPU（舉例H100）引入了與前代不同的新硬件單位和指示。提供硬件信息對Llama 3.1 70B或DeepSeek-V3的輸出影響不大。

真諦的是，他們發(fā)現(xiàn)OpenAI o1和DeepSeek-R1生成的部安分核使用了特定于硬件的指示和優(yōu)化。

R1在爽快50%的Level 1矩陣乘法問題中嘗試生成warp矩陣乘加（wmma）指示（下圖10），盡管大多數(shù)未能編譯。

在功能正確的生成中，R1和o1在每個級別產(chǎn)生了1-3個很是值，比Level 4基線快2倍以上。

總體而言，商酌團隊發(fā)現(xiàn)，與提供硬件信息比較，講話模子通過少樣本示例調(diào)遣戰(zhàn)術(shù)時闡明更好。

論斷

商酌東說念主員在本論文中提議了KernelBench，一個為講話模子驅(qū)動的內(nèi)核優(yōu)化奠定基礎(chǔ)的框架；他們評估了多種模子和門徑，分析了它們的上風(fēng)和局限性，并提供了創(chuàng)新契機的辦法。

總的來說，盡管大多數(shù)基準測試最終會達到充足，但KernelBench被設(shè)想為跟著新的AI責(zé)任負載的出現(xiàn)而動態(tài)演進。

他們的fastp主張不錯隨時辰調(diào)遣，以測量相干于日益先進的基線（即超出責(zé)任中使用的PyTorch基線）的加快閾值（p）。

由于PyTorch具備跨硬件平臺兼容性，KernelBench中基于PyTorch的任務(wù)不錯在每個新硬件平臺發(fā)布時進行評估。終末，與好多基準測試不同，在KernelBench上的見效徑直映射到坐褥價值和現(xiàn)實全國的影響（裁減資本并大范圍減少能耗）。

這些特點確保了KernelBench在約束演變的AI范圍中將保捏其價值。

下一步責(zé)任

商酌團隊示意在刻下可用模子下，KernelBench仍有權(quán)臣的創(chuàng)新空間。

最初，當年的責(zé)任不錯探索引誘先進的微和諧推理時代，包括智能體責(zé)任流（agentic workflows）。由于CUDA是一種低資源講話，當年責(zé)任開源更多高質(zhì)地數(shù)據(jù)將具有伏擊價值。

其次，在他們的實驗中，講話模子生成的是原始CUDA代碼。然則，當年的責(zé)任不錯探索使用其他編程籠統(tǒng)（舉例ThunderKittens、CUTLASS、Triton等）生成代碼是否能簡化生成問題，舉例使講話模子更容易欺詐張量中樞指示。

終末，商酌團隊的評估于今僅限于GPU，當年的責(zé)任不錯彭脹到其他硬件加快器。

作家先容

Anne Ouyang

Anne Ouyang面前是斯坦福大學(xué)籌劃機科學(xué)（CS）博士生，在Scaling Intelligence Lab（范圍化智能實驗室）進行商酌。

她的商酌酷好酷好主要匯集在可彭脹的自我創(chuàng)新機器學(xué)習(xí)系統(tǒng)，同期也平常珍重實證機器學(xué)習(xí)（empirical ML）和性能工程（performance engineering）。

她曾得回了MIT學(xué)士和碩士學(xué)位，并曾在NVIDIA cuDNN團隊責(zé)任，崇拜編寫CUDA內(nèi)核，用于加快GPU上的深度學(xué)習(xí)責(zé)任負載。

Simon Guo

Simon Guo是斯坦福大學(xué)籌劃機科學(xué)?？频囊荒挲X博士生，面前正在彭脹智能實驗室（Scaling Intelligence Lab）奴婢Azalia Mirhoseini西賓進行輪轉(zhuǎn)商酌。

他曾得回了UC伯克利電氣工程和籌劃機科學(xué)學(xué)士學(xué)位。他的商酌酷好酷好在籌劃機系統(tǒng)和機器學(xué)習(xí)。

最近，他在Cohere從事講話模子預(yù)進修責(zé)任。在此之前，他曾在蘋果公司設(shè)想GPU，在Anyscale引誘漫衍式系統(tǒng)體育游戲app平臺，并在NVIDIA DRIVE部門從事自動駕駛汽車的引誘責(zé)任。