嗨，我們要開(kāi)始了。我叫馬修·尼斯利。我是NVIDIA的深度學(xué)習(xí)compiler PM，今天我將介紹一些針對(duì)NVIDIA Tensorcores的使用方法。首先我要講一下Cutlass。我會(huì)給你一些背景和概述，為什么你可能會(huì)使用它，一些最新和即將推出的功能，然后我會(huì)概述一下開(kāi)放平臺(tái)Triton。如果你剛剛參加了上一場(chǎng)講座的話那你已經(jīng)是懂哥了。

好的，Cutlass是我們?cè)贕itHub上的開(kāi)源庫(kù)，用于在張量核上進(jìn)行編程。它最初于2017年發(fā)布，旨在改善Volta的可編程性。從那時(shí)起，它的應(yīng)用逐漸增長(zhǎng)。它已經(jīng)從一個(gè)供深度學(xué)習(xí)從業(yè)者使用的研究工具轉(zhuǎn)變?yōu)檎麄€(gè)生態(tài)系統(tǒng)中廣泛應(yīng)用的生產(chǎn)資產(chǎn)。

Cutlass由構(gòu)建模塊組成，可以根據(jù)您的需要使用gemm，卷積等，無(wú)論是從現(xiàn)成的還是自己設(shè)計(jì)的內(nèi)核。我們支持多種Epilogue模式以及在NVIDIA GPU上找到的所有數(shù)據(jù)計(jì)算類(lèi)型。我們最近發(fā)布了一個(gè)Python接口，我稍后會(huì)詳細(xì)介紹它。另外，我們還有一個(gè)性能分析器，你可以用它來(lái)找到最適合你使用情況的配置。Cutlass對(duì)于NVIDIA生態(tài)系統(tǒng)至關(guān)重要，你會(huì)在許多庫(kù)中找到它，例如cublas、CUSPARSE、cuTENSOR和DALI等等。那么，它是如何工作的呢？

Cutlass由五個(gè)抽象層構(gòu)成，其目的是最大限度地提高靈活性。首先，我們幾個(gè)月前發(fā)布了版本3，它有一個(gè)新的后端稱(chēng)為Cute，極大地簡(jiǎn)化了線程數(shù)據(jù)映射，并允許核心開(kāi)發(fā)人員專(zhuān)注于張量和算法的邏輯描述。（CuTe is a collection of C++ CUDA template abstractions for defining and operating on hierarchically multidimensional layouts of threads and data.）所以，分解這些抽象層，底層是原子，包含著PTX。tiled的MMA和copy；過(guò)了collective層，接下來(lái)是kernel層，在這里你可以將collective mainloop and a collective epilogue結(jié)合在一起。最后，在設(shè)備層面上，你會(huì)找到內(nèi)核配置、啟動(dòng)工具和設(shè)備層面保證的可移植性。

（譯者吐槽：這里如果不看官方文檔很多概念其實(shí)是看不懂的。。。。所以插播一下其他知識(shí)以便不會(huì)覺(jué)得太突兀。如果你不想看可以跳過(guò)，轉(zhuǎn)載自官方文檔 https://github.com/NVIDIA/cutlass/blob/main/media/docs/cutlass_3x_backwards_compatibility.md

Collective mainloop和epilogue是CUTLASS中用于執(zhí)行集合矩陣乘累加（MMA）操作的關(guān)鍵組件。

Collective mainloop（集合主循環(huán)）是一個(gè)循環(huán)結(jié)構(gòu)，用于在多個(gè)線程中執(zhí)行MMA操作。它負(fù)責(zé)將輸入矩陣切分成小塊，并在多個(gè)線程之間協(xié)調(diào)數(shù)據(jù)傳輸和計(jì)算操作。主循環(huán)使用MMA指令對(duì)這些小塊執(zhí)行矩陣乘累加操作，利用硬件的并行性和局部性來(lái)加速計(jì)算。主循環(huán)還處理線程同步和通信，以確保正確的數(shù)據(jù)依賴關(guān)系和結(jié)果的一致性。

Epilogue（收尾）是主循環(huán)之后的一系列操作，用于處理主循環(huán)的輸出結(jié)果。它可以執(zhí)行各種操作，如對(duì)結(jié)果進(jìn)行修正、縮放、舍入等。Epilogue的目的是將主循環(huán)的輸出轉(zhuǎn)換為最終的矩陣乘累加結(jié)果。CUTLASS提供了不同類(lèi)型的Epilogue，可以根據(jù)具體需求選擇適當(dāng)?shù)腅pilogue類(lèi)型。

通過(guò)將Collective mainloop和Epilogue組合在一起，CUTLASS能夠高效地執(zhí)行集合MMA操作，利用硬件的并行性和局部性來(lái)提高計(jì)算效率。這種組合的靈活性使得CUTLASS可以適應(yīng)不同的硬件架構(gòu)和應(yīng)用需求，并提供高性能的矩陣乘累加功能。

這里假設(shè)以一個(gè)gemm說(shuō)明為例：

// cutlass: ClusterTileM and ClusterTileN loops
//   are either rasterized by the hardware or scheduled by the kernel in persistent kernels.
// Parallelism over thread block clusters
for (int cluster_m = 0; cluster_m < GemmM; cluster_m += ClusterTileM) {
  for (int cluster_n = 0; cluster_n < GemmN; cluster_n += ClusterTileN) {

    // cutlass: mainloop that iterates over all k-tiles
    // No loop unrolling is performed at this stage
    for (int k_tile = 0; k_tile < size<2>(gmem_tensor_A); k_tile++) {

      // loops inside cute::gemm(tiled_mma, a, b, c); Dispatch 5: (V,M,K) x (V,N,K) => (V,M,N)
      // TiledMma uses the hardware instruction provided through its Mma_Atom
      // TiledMma's atom layout, value layout, and permutations define the iteration order
      for (int tiled_mma_k = 0; tiled_mma_k < size<2>(A); tiled_mma_k++) {
        for (int tiled_mma_m = 0; tiled_mma_m < size<1>(A); tiled_mma_m++) {
          for (int tiled_mma_n = 0; tiled_mma_n < size<1>(B); tiled_mma_n++) {

            // TiledMma's vector mode dispatches to the underlying instruction.
            mma.call(d, a, b, c);
          } // tiled_mma_n
        } // tiled_mma_m
      } // tiled_mma_k
    } // k_tile mainloop
  } // cluster_m
} // cluster_n

CUTLASS使用以下組件來(lái)表示上述循環(huán)嵌套（上面是一個(gè)gemm），這些組件針對(duì)數(shù)據(jù)類(lèi)型、布局和數(shù)學(xué)指令進(jìn)行了專(zhuān)門(mén)化。

在CUTLASS 3.0中，我們通過(guò)首先在內(nèi)核層面組合collective mainloop（cutlass: mainloop that iterates over all k-tiles）和collective epilogue ，然后使用主機(jī)端轉(zhuǎn)換器將它們包裝成一個(gè)GEMM內(nèi)核的handle 。

以下部分按照用戶實(shí)例化它們的順序描述了組裝一個(gè)內(nèi)核需要的組件，這個(gè)順序是：

組裝所需的collective mainloop and epilogues。

將它們組合在一起構(gòu)建一個(gè)內(nèi)核類(lèi)型。

使用設(shè)備層轉(zhuǎn)換器包裝內(nèi)核。

Collective是“mma atoms和copy atoms 被切分到的最大線程集合”。也就是說(shuō)，它是一個(gè)最大數(shù)量的線程網(wǎng)格，通過(guò)利用硬件特性進(jìn)行加速通信和同步來(lái)進(jìn)行協(xié)作。這些硬件特性包括：

異步數(shù)組復(fù)制（例如，從全局內(nèi)存到共享內(nèi)存）；

用于位于共享內(nèi)存中的小塊的MMA指令；

用于集群、線程塊和/或warp的同步操作；和/或

硬件加速（例如屏障），以確保滿足異步操作之間的數(shù)據(jù)依賴關(guān)系。

Collective使用TiledMma和TiledCopy API（來(lái)訪問(wèn)在塊上執(zhí)行復(fù)制和MMA操作。

cutlass::CollectiveMma類(lèi)是集合矩陣乘累加（MMA）mainloop的主要接口。這里的“主循環(huán)”指的是在偽代碼中靠近本文頂部的“cluster tile k”循環(huán)。算法可能需要對(duì)多個(gè)塊進(jìn)行循環(huán)的情況會(huì)在這里發(fā)生。

更多請(qǐng)自行查看文檔

為什么要使用Cutlass呢？這可能是最常見(jiàn)的問(wèn)題。cublas將擁有最佳的開(kāi)箱體驗(yàn)。它將有更快的上市時(shí)間。它在不同架構(gòu)之間提供了可移植性保證。它有一組基于您的參數(shù)選擇最佳內(nèi)核的啟發(fā)式算法。所以我告訴很多客戶的是，如果cublas能滿足您的需求，就使用它。

（譯者：以防看不懂放上GPT的解釋?zhuān)?/p>

CUTLASS和CUBLAS是兩個(gè)用于在NVIDIA GPU上進(jìn)行矩陣運(yùn)算的庫(kù)，它們有以下區(qū)別：

開(kāi)發(fā)者：CUTLASS是由NVIDIA開(kāi)發(fā)和維護(hù)的開(kāi)源項(xiàng)目，而CUBLAS是NVIDIA官方提供的閉源庫(kù)。

靈活性和可配置性：CUTLASS提供了更高級(jí)別的靈活性和可配置性，允許用戶自定義和優(yōu)化矩陣運(yùn)算的細(xì)節(jié)。它提供了底層的矩陣運(yùn)算原語(yǔ)和算法的實(shí)現(xiàn)，使用戶可以根據(jù)特定需求進(jìn)行定制和優(yōu)化。CUBLAS則提供了更高級(jí)別的抽象和易用性，適用于常見(jiàn)的矩陣運(yùn)算任務(wù)。

性能優(yōu)化：CUTLASS注重性能優(yōu)化和硬件特性的利用。它提供了更多的配置選項(xiàng)和優(yōu)化策略，使用戶能夠根據(jù)具體的硬件架構(gòu)和應(yīng)用需求進(jìn)行性能優(yōu)化。CUTLASS還提供了針對(duì)深度學(xué)習(xí)任務(wù)的特殊優(yōu)化，如半精度浮點(diǎn)計(jì)算（FP16）和Tensor Core加速。CUBLAS也進(jìn)行了一些性能優(yōu)化，但它更注重提供易用性和通用性。

支持的功能：CUTLASS提供了更多的功能和算法選項(xiàng)，包括矩陣乘累加（MMA）、卷積等。CUBLAS則提供了一組預(yù)定義的矩陣運(yùn)算函數(shù)，如矩陣乘法、矩陣向量乘法等。

開(kāi)源性：CUTLASS是開(kāi)源的，用戶可以訪問(wèn)其源代碼并參與社區(qū)貢獻(xiàn)和討論。CUBLAS是閉源的，用戶無(wú)法訪問(wèn)其底層實(shí)現(xiàn)。)

如果您需要最大的靈活性，比如自定義epilogue，在cublas中并不存在，那么就使用Cutlass。雖然它需要花費(fèi)一些時(shí)間來(lái)啟動(dòng)和運(yùn)行，但您可以對(duì)數(shù)據(jù)傳輸和操作擁有最大的控制權(quán)。

在PyTorch生態(tài)系統(tǒng)中，你在哪里可以找到Cutlass呢？在高層級(jí)上，你會(huì)在eager模式下找到一些稠密和稀疏操作，并且目前有一個(gè)PR正在將Cutlass作為Inductor的另一種后端添加進(jìn)去。AItemplate（meta的torch backend codegen工具）在開(kāi)發(fā)過(guò)程中使用了每一層的一些特性。Xformer的內(nèi)存高效注意力是在Cutlass上開(kāi)發(fā)的。最后，PyTorch geometric是我們group Gemm的早期采用者之一。

我之前提到了Python接口。Cutlass的最大痛點(diǎn)之一是C++模板。通過(guò)我們的Python接口，我們大大減少了開(kāi)始所需的默認(rèn)值。這是一個(gè)基本的gemm示例，你可以看到它可能是所需參數(shù)的三分之一多。

using Gemm = typename gemm::DefaultGemmUniversal<
    half, layout::RowMajor, ComplexTransform::kNone,
    half, layout::RowMajor, ComplexTransform::kNone,
    half, layout::RowMajor,
    arch::OpClassTensorOp, arch::Sm8e,
    gemm::GemmShape<256, 128, 64>,
    gemm::GemmShape<64, 64, 64>,
    gemm::GemmShape<16, 8, 16>,
    epilogue::LinearCombination
    gemm::GemmIdentityThreadblockSwizzle<1>,
    arch::OpMultiplyAdd>::GemmKernel;

# 創(chuàng)建GEMM plan
plan = cutlass.op.Gemm(element=torch.float16, layout=cutlass.LayoutType.RowMajor)

# 更改swizzling functor
plan.swizzling_functor = cutlass.swizzle.ThreadblockSwizzlestreamk

# 添加fused ReLU
plan.activation = cutlass.epilogue.relu

Python接口的另一個(gè)目標(biāo)是改善與PyTorch等框架的集成。可以通過(guò)新的Cutlass emit PyTorch方法來(lái)實(shí)現(xiàn)。右側(cè)您將看到可以使用Python接口聲明PyTorch中的group gemm并給出PyTorch CUDA擴(kuò)展。

之后我們將負(fù)責(zé)生成源代碼，為PyTorch擴(kuò)展提供輸入，并提供一個(gè)腳本來(lái)生成該P(yáng)yTorch擴(kuò)展。

在Cutlass中最新的功能是我們稱(chēng)為epilogueVisitor Tree。這將允許用戶使用基本的epilogue單元來(lái)開(kāi)發(fā)復(fù)雜的epilogue。它是一組小的計(jì)算、加載和存儲(chǔ)操作，可以生成常見(jiàn)或自定義的epilogue。它在Ampere和Hopper架構(gòu)上和C++和Python接口中已經(jīng)可以使用。我們?cè)诂F(xiàn)有的配置中，利用峰值效果非常好。下面是一個(gè)示例，在最新版本的Cutlass 3.2和Cuda Toolkit 12.2以及H100上使用性能分析器和默認(rèn)參數(shù)。

您可以看到對(duì)于所有這些用例，我們大約達(dá)到了90%的峰值利用率。我們還努力確保沒(méi)有性能下降，并會(huì)定期發(fā)布優(yōu)化版本。

我們的下一個(gè)發(fā)布版本是3.3。3.3中最重要的功能是我們稱(chēng)之為混合輸入gemm。這是一個(gè)常見(jiàn)的需求，在這個(gè)功能下，你可以為A和B矩陣使用不同的數(shù)據(jù)類(lèi)型。例如，A可以是FP16。B可以是int8，而在gemm中，我們將向上轉(zhuǎn)換到BF16的操作。我們還對(duì)我們的lower alignment gemm進(jìn)行了性能改進(jìn)。然后在今年12月，我們將推出grouped 和 ptr-array gemm，以及為Hopper優(yōu)化的四個(gè)性能卷積。明年年初，我們將為Ada提供FP8支持。我們將進(jìn)行卷積的權(quán)重和偏差（W&D grad）優(yōu)化，支持稀疏數(shù)據(jù)。我們還有一個(gè)稱(chēng)為portablepipeline的新功能。portablepipeline是我們?yōu)橄Ｍ诩軜?gòu)上實(shí)現(xiàn)可移植性的用戶提出的建議。在GTC Talk上會(huì)有更多關(guān)于這個(gè)功能的信息。最后，在明年第二季度，Cutlass開(kāi)發(fā)者和初學(xué)者需要的是更好的文檔。我們會(huì)進(jìn)行全面更新，這將涵蓋C++和Python接口。

OpenAI Triton是一個(gè)令人興奮的新的類(lèi)似于Python的編程語(yǔ)言，供用戶開(kāi)發(fā)在Python中針對(duì)NVIDIA TensorCores的內(nèi)核。這樣一來(lái)，開(kāi)發(fā)者可以專(zhuān)注于更高級(jí)別的邏輯。OpenAI Triton Compiler會(huì)處理許多性能優(yōu)化，讓開(kāi)發(fā)者不用操心。NV和openai在大力合作推進(jìn)一切中。

到此結(jié)束，謝謝。