剪輯：剪輯部
斯坦福和普林斯頓商酌者發現，DeepSeek-R1生成的自界說CUDA內核，完爆了o1和Claude 3.5 Sonnet，拿下總排行第一。誠然面前只可在不到20%任務上卓絕PyTorch Eager基線，但GPU編程加快自動化的按鈕，照舊被按下！
近日，來自斯坦福和普林斯頓的商酌者發現，DeepSeek-R1照舊能生成自界說CUDA內核了，而且還在一眾推理模子中，徑直拿下了TOP 1！
緊隨后來，OpenAI o1和Claude 3.5 Sonnet鑒識排第二和第三。
具體過程，便是給定一個PyTorch法子，讓模子對其優化，然青年景一個包含自界說CUDA內核的PyTorch版塊。
在此期間中，模子不錯開脫決定優化哪些操作，以提高籌劃效率。
輔導模子生成GPU內核，后勁紛亂
如今，傳統東說念主工優化內核的格式，在效率上照舊不及以應酬多數暴露的AI模子架構和硬件平臺。
既然是為了LLM進行優化，那么，咱們是否也冒失借助LLM來模擬AI工程師的責任經過，憑借編譯器反應、硬件規格等豐富的信息，自動編寫出準確且經過優化的內核代碼呢？
為此，商酌團隊提議了一種全新的KernelBench框架，用于生成和評估不同AI任務（單個操作、操作序列、端到端架構）的內核，并模擬了工程師迭代優化的過程。
論文地址：https://arxiv.org/abs/2502.10517
GPU的本體，是硬件依賴的。因此，商酌者們但愿嘗試，看是否能通過以下格式，輔導模子生成GPU內核。
最初，向模子提供硬件信息（如內存帶寬、TFLOPS），以針對特定GPU（A100、H100等）進行優化。
然后，要讓模子在崎嶇文中展示代表性的內核優化手段，舉例矩陣乘法中的分塊（tiling）或Flash Attention中的在線softmax。
商酌者們發現，唯一更強勁的模子，會偶爾闡明出欺詐這些優化的能力。
崎嶇滑動稽察
比如，DeepSeek-R1未必會使用特定于硬件的指示（如Tensor Core的wmma），但時常無法正確編譯或使用它們，從而脅制了最終性能。
總的來說，商酌發現，前沿模子在KernelBench上的開箱即用性較差，OpenAI o1和DeepSeek-R1在不到20%的任務上高出PyTorch Eager基線。
這些模子生成的內核存在多數引申失實、功能正確性問題，何況無法進行特定平臺的優化。
具體來說，商酌者發現：
對模子而言，編寫功能正確的內核仍然具有挑戰性；
模子通過優化展示了生成高性能內核的后勁；
欺詐反應關于減少引申失實和發現更快的決策很伏擊。
天然，KernelBench面前還只是讓GPU加快奔走的發軔，但亦然讓通盤GPU編程自動化的肇始催化劑。
令東說念主歡快的是，面前照舊有了好多新的責任，專注于處理KernelBench中觸及的問題。
比如2月12日，英偉達就發出博客著述，探討奈何使用DeepSeek-R1進行GPU內核自動生成與推理時scaling。
隨后在2月12日，Meta也發文測試了前沿模子編寫GPU內核方面的性能，他們發現，最好模子不錯在KernelBench上提供平均1.8倍的加快。
Sakana AI更是推出「AI CUDA工程師」，讓AI我方寫代碼優化CUDA內核，速率宣稱比PyTorch原生完結快了10-100倍。
如棋布星陳般出現的商酌標明，如今，咱們照舊參加了AI驅動加快AI的新紀元！
在當年，KernelBench還將捏續演進。它不會僅限于刻下網羅的250個問題，還不錯彭脹到新的AI任務。與此同期，評測主張fast_p也不錯跟著時辰的推移進行調遣，提高加快門檻，以鼓勵更高效的優化決策
KernelBench：AI內核生成框架
KernelBench是一個開源框架，旨在評估LLM在編寫GPU內核方面的能力。
任務格式
KernelBench包含250個任務，涵蓋了多樣AI責任負載，何況易于彭脹到新的責任負載。
下圖1展示了KernelBench評估講話模子（LM）生成高性能GPU內核的能力。KernelBench條件講話模子為給定的標的PyTorch模子架構生成優化的CUDA內核，并進行自動化評估。
· 任務輸入
給定一個AI責任負載，任務的輸入是用PyTorch編寫的參考完結。效法商酌東說念主員的責任經過，PyTorch代碼包含一個秉承自torch.nn.Module ()的名為Model的類，其中圭臬的__init__和 forward () 函數（以及任何輔助函數）被填充為AI責任負載的PyTorch操作。
AI算法閑居在大型張量數據上進行操作。責任負載的最優內核取決于張量的大小和數據類型（如BF16、FP8）。因此，每個任務還包含get_inputs ()和get_init_inputs ()函數，用于指定內核需要處理的精準輸入張量。
· 任務輸出
給定輸入，LLM需要輸出一個秉承自torch.nn.Module ()的名為ModelNew的新類，其中包含自界說優化。舉例，LLM不錯在forward ()函數中使用PyTorch的CUDA-C彭脹來集成內聯內核調用。
為了見效完成任務，模子需要服氣（1）Model類中的哪些操作最能從優化中受益；（2）奈何優化這些操作。LLM不錯使用任何硬件高效時代（如交融和分塊）或專用指示（如張量中樞）以及任何編程庫（如PTX、CUDA、CUTLASS、Triton、ThunderKittens）。
崎嶇滑動稽察
任務遴薦
這些任務左證包含的基本操作或PyTorch庫函數的數目分為三個級別。
Level 1包含100個單個基本操作，如卷積、矩陣乘法等AI基礎構建塊。誠然PyTorch調用了經過優化的閉源內核，讓LLM卓絕基線具有挑戰性，但如果能生成開源內核，將有伏擊價值。
Level 2包含100個操作序列，如卷積、ReLU和Bias的組合，這些操作不錯交融成一個內核以提高性能。
由于基于編譯器的器具（如PyTorch編譯器）在交融方面相稱靈驗，LLM要在這方面卓絕它們也具有挑戰性。然則，LLM可能會提議更復雜的算法。
Level 3包含50個竣工的機器學習架構，如AlexNet和MiniGPT等，這些架構在運行進修和推理時對內核的性能條件極高。
評估主張
KernelBench引入了一個新的評估主張fast_p，估量生成的內核中功能正確且加快大于閾值p的任務比例。
通過調遣閾值參數p，商酌者不錯評估不同加快閾值下的內核性能，并捕捉加快漫衍。
fast_0終點于LLM的正確率，它估量代碼功能正確的任務比例，而不研討其速率。在試驗評估中，閑居以p=1動作發軔。
LLM在KernelBench上的闡明
商酌東說念主員對一系列LLM在KernelBench上進行了評估，收尾自滿，面前的LLM在生成正確且優于PyTorch基線速率的內核方面仍有艱苦。
在一次性基線評估中，LLM生成的內核平均在不到20%的任務中比PyTorch Eager更快。這標明，僅靠膚淺領導，LLM很難在性能上卓絕傳統的PyTorch內核。
LLM生成的內核存在多數的引申失實和功能正確性問題，經常由于膚淺的編譯器和運行時失實而失敗。
引申失實包括CUDA/nvcc/Python編譯時失實、CUDA內存違法和運行時失實等；正確性失實則主要闡明為輸出張量局勢和值不匹配。
推理模子（o1，R1）生成的失實處理決策（<55%）比其他模子（>70%）少。然則，這主淌若由于引申失敗的情況較少。在功能正確性方面，總計LLM王人瀕臨相似進度的艱苦。
在性能方面，模子生見遵守正確的內核在多數情況下也未能優于PyTorch基線。
跟著p的提高，模子生成的內核中能達到條件的比例越來越低。在p=1時，在總計KernelBench級別中，不到15%的LLM生成內核優于PyTorch。
推理模子閑居在提供加快方面優于其他LLM，但總體仍有不及。
模子生成的內核在不同硬件平臺上的通用性欠安。DeepSeek-R1生成的內核在NVIDIA L40S上完結了36%的加快，而在NVIDIA A10G上則為47%。
這標明LLM在生成特定標的硬件的高效內核方面還存在很大的提拔空間。
模子能力分析
測試時欺詐KernelBench環境反應
正如上頭不雅察到的，引申失敗是LM生成的內核中最常見的失敗模式。
KernelBench提供的環境允許網羅豐富的信號，包括編譯器失實、正確性查驗和運行時性能分析主張，總計這些王人不錯反應給LM以匡助它處理內核故障。
為了探索LM奈何欺詐這些反應，商酌團隊評估和比較了兩個基線：第一個令LM為每個KernelBench任務生成多個并行樣本，另一個通過允許LM欺詐引申反應漸漸創新，挨次為每個KernelBench任務生成內核。
· 重迭采樣
KernelBench環境營救對LM生成的內核進行法子化考據，允許商酌團隊網羅和評估每個任務的多個LM生成。他們使用fastp@k評估這種重迭采樣門徑。重迭采樣有助于LM發現更多快速且正確的處理決策。
如下圖4所示，跟著k值的加多，在DeepSeek-V3和Llama 3.1 70B的三個級別上，通過高溫度參數重迭采樣不錯提拔fast1的性能。
值得珍視的是，在Level 2上，DeepSeek-V3在k=100個樣本時達到了37%的fast1，而在單次領導基線中僅為4%。
通過查驗樣本，咱們發現高溫采樣有助于探索處理決策空間，加多了生成具有更好優化的無失實內核的契機。然則，如果一個模子處理任務的固有概率相稱低，只是加多采樣預算的影響有限。
舉例，即使嘗試了100個樣本，DeepSeek-V3也從未冒失為Level 1中的一組34個卷積變體生成任何正確的處理決策。
· 生成收尾的迭代優化
KernelBench環境相稱合乎網羅編譯器反應、引申失實和使用PyTorch分析器等器具進行的時辰分析，動作果然信號（ground-truth signals）。
商酌東說念主員商酌了欺詐這些反應是否能匡助講話模子（LMs）迭代優化其生成收尾。
下圖5自滿，KernelBench框架使模子冒失在迭代優化過程中領受并欺詐反應。這些果然信號包括NVCC編譯器失實信息、引申統計數據（舉例正確性查驗和掛鐘時辰），以及PyTorch分析器（操作時辰認識）。
他們在多輪過程中為模子提供每次生成的反應：在運轉生成后，向模子提供其之前的生成收尾G，以及刻下生成對應的編譯器/引申反應E和/或分析器輸出P。
然后將每次生成偏激后續反應界說為一輪（turn），并在N輪內運行這一迭代優化過程。欺詐引申反應有助于減少失實，并隨時辰提拔舉座加快效率。
商酌東說念主員不才表2中查驗了第N=10輪時的fast1行為，發現迭代優化在不同模子和KernelBench的各個級別上均捏續提拔了性能。
DeepSeek-R1在Level 2上的創新最為權臣，其中引申反應E和分析器反應P的組合將fast1從36%提拔至72%（如下圖6所示）。
此外，通過分析迭代優化軌跡，他們發現模子在引申反應E的匡助下能更靈驗地自我矯正，尤其是在建設與引申失實籌商的問題上。
DeepSeek-R1在Level 1和Level 2上，經過10輪優化后，能在高出90%的任務中生見遵守正確的內核（下表9）。
然則，剩余的失實內核簡直老是由于功能不正確而失敗，這可能是因為正確性反應的顆粒度不如引申失敗信息考究。
· 比較重迭采樣與迭代優化
在上表2中，商酌東說念主員比較了在固定10次推理調用預算下重迭采樣和迭代優化的效率。兩種門徑相較于單次基線（one-shot baseline）均取得了權臣創新，其中迭代優化在6個案例中的5個中闡明更優。
然則，他們最終發現，測試時門徑的效率本體上依賴于基礎模子的質地。
舉例，在重迭采樣中，DeepSeek-V3在總計三個級別上恒久優于Llama-3.1 70B。相似地，在迭代優化中，DeepSeek-R1通過反應E和P捏續創新，而DeepSeek-V3和Llama-3.1 70B并非總能從這些信息中獲益。
提供硬件學問生成硬件高效內核
昭彰，講話模子在生成硬件高效內核方面闡明存限。
這可能是由于進修數據中內核代碼的稀缺性，以及最好內核可能需要左證硬件平臺的特定屬性而變化。
在本案例商酌中，商酌團隊探索了提供以下內容的效率：（1）提供內核工程最好實踐的示例，并將其置于（講話模子的）崎嶇文之中；（2）提供詳備的硬件規格施展，并將其置于（講話模子的）崎嶇文之中。
· 硬件感知的崎嶇文示例
編寫細密的內核閑居使用交融（fusion）、分塊（tiling）、重籌劃（recompute）和異步（asynchrony）等時代來最大化性能。
具體來說，商酌東說念主員納入了三個崎嶇文示例：使用操作交融的GeLU、使用分塊的矩陣乘法，以及展示分享內存I/O經管的最小Flash-Attention內核。
收尾則自滿，崎嶇文示例裁減了講話模子的舉座fast1分數，因為模子嘗試了更激進的優化戰術，但導致更多引申失敗。與上頭基線生成的代碼比較，OpenAI o1在使用少樣本示例時生成的代碼平均長度加多了25%。
然則，在正確的處理決策中，講話模子應用了一些真諦的優化：他們發現，在KernelBench Level 1的77%的GEMM變體中，o1應用了分塊并提拔了速率，優于單次基線。在Level 2，o1在11個問題上應用了激進的分享內存I/O經管，并冒失卓絕PyTorch Eager。
· 指定硬件信息
正如上頭籌謀過的，內核性能因硬件平臺而異。
舉例，FlashAttention-2從NVIDIA A100搬動到H100 GPU時硬件欺詐率下落了47%。FlashAttention-3是一個全王人不同的算法，專為H100編寫。
在本商酌中，商酌團隊探討講話模子是否能欺詐崎嶇文中的以下信息生成創新的內核：（1）硬件規格，舉例GPU類型（H100、A100等）、內存大小、帶寬、TFLOPS；（2）硬件學問（舉例線程、線程束、線程塊、流多處理器的界說）。
收尾自滿，模子很少生成針對底層硬件優化的內核，這標明當年模子仍有創新空間。
某些新一代GPU（舉例H100）引入了與前代不同的新硬件單位和指示。提供硬件信息對Llama 3.1 70B或DeepSeek-V3的輸出影響不大。
真諦的是，他們發現OpenAI o1和DeepSeek-R1生成的部安分核使用了特定于硬件的指示和優化。
R1在爽快50%的Level 1矩陣乘法問題中嘗試生成warp矩陣乘加（wmma）指示（下圖10），盡管大多數未能編譯。
在功能正確的生成中，R1和o1在每個級別產生了1-3個很是值，比Level 4基線快2倍以上。
總體而言，商酌團隊發現，與提供硬件信息比較，講話模子通過少樣本示例調遣戰術時闡明更好。
論斷
商酌東說念主員在本論文中提議了KernelBench，一個為講話模子驅動的內核優化奠定基礎的框架；他們評估了多種模子和門徑，分析了它們的上風和局限性，并提供了創新契機的辦法。
總的來說，盡管大多數基準測試最終會達到充足，但KernelBench被設想為跟著新的AI責任負載的出現而動態演進。
他們的fastp主張不錯隨時辰調遣，以測量相干于日益先進的基線（即超出責任中使用的PyTorch基線）的加快閾值（p）。
由于PyTorch具備跨硬件平臺兼容性，KernelBench中基于PyTorch的任務不錯在每個新硬件平臺發布時進行評估。終末，與好多基準測試不同，在KernelBench上的見效徑直映射到坐褥價值和現實全國的影響（裁減資本并大范圍減少能耗）。
這些特點確保了KernelBench在約束演變的AI范圍中將保捏其價值。
下一步責任
商酌團隊示意在刻下可用模子下，KernelBench仍有權臣的創新空間。
最初，當年的責任不錯探索引誘先進的微和諧推理時代，包括智能體責任流（agentic workflows）。由于CUDA是一種低資源講話，當年責任開源更多高質地數據將具有伏擊價值。
其次，在他們的實驗中，講話模子生成的是原始CUDA代碼。然則，當年的責任不錯探索使用其他編程籠統（舉例ThunderKittens、CUTLASS、Triton等）生成代碼是否能簡化生成問題，舉例使講話模子更容易欺詐張量中樞指示。
終末，商酌團隊的評估于今僅限于GPU，當年的責任不錯彭脹到其他硬件加快器。
作家先容
Anne Ouyang
Anne Ouyang面前是斯坦福大學籌劃機科學（CS）博士生，在Scaling Intelligence Lab（范圍化智能實驗室）進行商酌。
她的商酌酷好酷好主要匯集在可彭脹的自我創新機器學習系統，同期也平常珍重實證機器學習（empirical ML）和性能工程（performance engineering）。
她曾得回了MIT學士和碩士學位，并曾在NVIDIA cuDNN團隊責任，崇拜編寫CUDA內核，用于加快GPU上的深度學習責任負載。
Simon Guo
Simon Guo是斯坦福大學籌劃機科學專科的一年齒博士生，面前正在彭脹智能實驗室（Scaling Intelligence Lab）奴婢Azalia Mirhoseini西賓進行輪轉商酌。
他曾得回了UC伯克利電氣工程和籌劃機科學學士學位。他的商酌酷好酷好在籌劃機系統和機器學習。
最近，他在Cohere從事講話模子預進修責任。在此之前，他曾在蘋果公司設想GPU，在Anyscale引誘漫衍式系統體育游戲app平臺，并在NVIDIA DRIVE部門從事自動駕駛汽車的引誘責任。