1 GPU基礎

1.1 GPU與CPU在數據處理中的對比

如今，各個領域——科學、商業(yè)、工程、媒體，甚至我們日常的網頁瀏覽——都依賴于不斷增長的數據量。單個桌面或服務器 CPU 就能處理整個工作流程的時代正在消逝。因此，隨著數據量從兆字節(jié)（MB）到千兆字節(jié)（GB），再到兆兆字節(jié)（TB），我們會發(fā)現我們的腳本和應用程序正難以跟上這種增長的步伐。這不僅僅關乎數據量；還關乎數據的復雜性、速度以及快速獲得結果的重要性。

我們的任務可能涉及處理醫(yī)學影像、轉換大型數據集進行分析、訓練或部署神經網絡、運行模擬，或向數百萬人傳輸流媒體視頻。即使是多核 CPU，其設計初衷也并非為了處理如此大量的并行重復性工作。您可能已經開始更多地了解 GPU 加速，以及它如何釋放近乎神奇的性能，將原本需要數小時才能完成的任務縮短到幾分鐘甚至幾秒鐘。

在開始編寫代碼之前，我們應該了解為什么如此多的應用程序正在從 CPU 轉向 GPU。本主題將幫助我們理解主要驅動力、硬件差異以及它們對我們日常數據挑戰(zhàn)的實際影響。

1.1.1 CPU 和 GPU 如何處理數據？

計算機系統(tǒng)的中央處理器 (CPU) 通常被稱為設備的“大腦”。它功能極其豐富，能夠進行復雜的決策，運行整個操作系統(tǒng)，并同時管理數十個軟件線程。CPU 內核功能強大且靈活。它配備了復雜的控制邏輯、巨大的緩存和極高的時鐘頻率。如今，臺式機和服務器中的大多數 CPU 都擁有 4 到 64 個內核，每個內核都可以處理獨立的指令流。

圖片來源：https://developer.nvidia.com/blog/cuda-refresher-reviewing-the-origins-of-gpu-computing

然而，CPU 針對順序處理進行了優(yōu)化。您可能會發(fā)現我們的代碼大量時間都花在循環(huán)或單線程例程上，而這些正是 CPU 的優(yōu)勢所在。但是，當處理海量數組時（例如過濾十億個圖像像素或對大型矩陣進行乘法運算），CPU 的設計就暴露出其局限性。雖然速度很快，但每個核心每秒只能處理有限的指令。添加更多核心固然有幫助，但很快就會達到收益遞減的臨界點。

另一方面，GPU（圖形處理單元）采用了一種截然不同的方法。因此，它不再擁有幾個“智能”核心，而是擁有數千個設計用于協(xié)同工作的簡單、輕量級核心。雖然每個 GPU 核心的性能不如單個 CPU 核心強大，但它們擅長同時對海量數據塊執(zhí)行相同的操作。這被稱為單指令多數據 (SIMD) 并行。

1.1.2 CPU 擴展的限制

有人可能會問：為什么我們不能直接添加更多 CPU 核心呢？畢竟，如今的服務器 CPU 都配備了數十個核心，而云服務提供商提供的服務器則配備了數百個虛擬 CPU (vCPU)。然而，問題遠比簡單的“越多越好”要復雜得多。

首先，CPU 的擴展成本很高。每個核心都需要復雜的控制單元和大型低延遲緩存，這會占用大量的硅片空間并消耗大量電力。其次，隨著核心數量的增加，我們會遇到阿姆達爾定律——程序的速度取決于其最慢、最連續(xù)的部分。除非代碼的每個部分都能獨立運行，否則總會存在阻礙系統(tǒng)完美擴展的瓶頸。

例如，我們可能使用線程或多處理等庫編寫了多線程 Python 腳本。隨著線程數量的增加，我們很快就會注意到全局鎖、線程爭用以及臭名昭著的全局解釋器鎖 (GIL) 限制了這些優(yōu)勢。在線程之間傳輸數據的過程會產生開銷，從而減少可用的內存帶寬。實際上，CPU 可以處理數十個并行線程，但無法處理數千個。

1.1.3 GPU 勝過 CPU

這正是 GPU 發(fā)揮作用的地方。想象一下，處理一幅 4,000 x 4,000 像素的圖像。如果我們要對每個像素應用變換，則需要進行 1600 萬次獨立計算。這可能包括調整亮度或應用濾鏡。即使擁有 16 或 32 個核心，CPU 也必須將這些工作負載劃分到少數幾個線程上。管理如此多的小任務并在核心之間移動數據的開銷將是巨大的。
然而，GPU 正是為這項任務而生。憑借數千個計算核心，GPU 可以將每個像素的操作分配給各自的線程。突然之間，整個圖像就可以并行處理，就像每個像素都有自己的處理器一樣。數字上的區(qū)別顯而易見：CPU 需要花費數秒才能完成的工作負載，在 GPU 上只需幾分之一秒即可完成。這種架構不僅適用于圖形處理。GPU 的對于任何遵循“相同操作，多個數據點”范式的工作負載來說，并行能力都是一大福音。這包括科學計算、神經網絡推理、統(tǒng)計模擬和大規(guī)模數據分析。

1.1.4 深入了解硬件

為了充分利用 GPU 編程，了解一些硬件內部結構會有所幫助。現代 CPU 可能擁有 8-32 個核心，每個核心的運行頻率為 3-4 GHz，并擁有復雜的流水線和數兆字節(jié)的緩存。相比之下，典型的 NVIDIA GPU 可能擁有 5,000 到 10,000 個簡單的 CUDA 核心，以較低的時鐘速度運行，但被分組到流多處理器 (SM) 中。每個 SM 調度和管理 32 個線程組（稱為 Warp），硬件可以同時保持多個 Warp 處于“運行”狀態(tài)。

GPU 通過 PCI Express 連接到系統(tǒng)的其余部分，雖然速度很快，但仍然比片上內存帶寬慢得多。您可以在 GPU 上訪問多種類型的內存：全局內存（大容量、高延遲、高帶寬）、共享內存（小容量、低延遲，位于 SM 內）以及寄存器（超高速，每個線程）。硬件旨在最大化吞吐量；即通過隱藏內存延遲并將計算與數據移動重疊來實現每秒處理的數據總量。

1.1.5 為什么 GPU 編程現在如此重要？

如果我們主要使用 Python 工作，我們就會知道使用 NumPy 或 Pandas 等庫操作數據是多么容易。隨著數據的增長，我們發(fā)現 NumPy 操作會變得緩慢，尤其是在執(zhí)行大型矩陣乘法、約簡或對大型數組進行元素級計算等操作時。GPU 編程曾經只為那些愿意深入研究 C++ 和復雜 CUDA API 的人所準備。值得慶幸的是，像 CuPy 和 PyCUDA 這樣的庫現在將 GPU 的強大功能直接引入我們的 Python 代碼中，讓我們能夠使用熟悉的語法進行并行計算。您通常可以采用 NumPy 編寫的代碼，替換掉 import 語句，即可看到顯著的加速——所有這些都無需離開 Python 的舒適區(qū)。

• 機器學習與深度學習推理
  假設我們正在部署一個神經網絡，每秒處理數千張圖像進行對象檢測或分類。深度神經網絡的前向傳播包含許多大型矩陣乘法和元素激活。如果沒有大規(guī)模集群，CPU 就無法在這種規(guī)模下進行實時推理。專為并行浮點運算而設計的 GPU 可以在極短的時間內提供結果，支持從云端到邊緣的大規(guī)模 AI 服務。

• 高性能數據分析
  您可以處理 ETL 管道、處理日志或對數百萬條記錄進行統(tǒng)計分析。分組、聚合、過濾和直方圖計算等任務與 GPU 的優(yōu)勢完美契合。RAPIDS 和 CuDF 等 Python 庫利用 GPU 進行數據庫式操作，使我們能夠加速分析工作負載，使其遠遠超出 CPU 的能力。

• 實時可視化與渲染
  交互式可視化、3D 渲染和游戲都依賴于 GPU 每秒數十次更新數百萬像素的能力。即使在圖形處理領域之外，科學家和工程師也使用 GPU 來可視化模擬結果、繪制大型點云或制作分子動力學動畫。CPU 根本無法應對這些高要求、高度并行的工作負載。

• 科學計算與模擬
  物理模擬——天氣模型、流體動力學、蒙特卡洛模擬——需要更新數百萬個粒子或網格點的狀態(tài)。GPU 在這些高度并行的問題上表現出色，讓我們能夠運行更精確或更高分辨率的模型，而無需等待數小時才能獲得結果。

在 GPU 上運行代碼時，并非總能獲得 1000 倍的加速。關鍵在于將工作負載與硬件的優(yōu)勢相匹配。對于足夠大的數組，我們可以預期向量加法、矩陣乘法和約簡等常見任務的加速可達 10 倍到 100 倍。對于較小的數據，將數據移入和移出 GPU 的開銷可能會抵消任何優(yōu)勢。

為了親眼見證這一點，我們將很快用 Python 運行一個基準測試，比較 CPU（NumPy）和 GPU（CuPy）上簡單的數組操作。通過繪制運行時間與數組大小的關系圖，我們發(fā)現 GPU 在處理大數據時絕對“勝出”，而對于傳輸開銷占主導地位的小任務，CPU 有時速度更快。

1.1.6 GPU 編程有何不同？

GPU 擅長處理具有以下特點的工作負載：
● 多個數據點的操作相同（數據并行）
● 每個計算都是獨立的（幾乎沒有線程間通信）
● 您需要處理大型數據集，而不僅僅是少數幾個項目

GPU 不擅長的地方：本質上是順序的、嚴重依賴復雜分支的或需要線程間持續(xù)通信的工作負載。對于這些工作負載，CPU 仍然至關重要。
隨著我們轉向 GPU 編程，我們的思維方式發(fā)生了轉變。你不再考慮“一次一個操作”，而是問，“我能不能用一種方式來寫，讓所有元素都“一起用嗎？”你開始關心內存布局、合并以及如何構建工作以最小化等待。在 Python 中，我們可以使用 CuPy 或 PyCUDA 自動完成大部分工作，但當我們了解如何將算法與硬件對齊時，我們將看到更好的結果。
我們經常會想到一些常見問題和真正的擔憂，例如：
● GPU 編程總是有用嗎？
● 學習 GPU 編程很難嗎？
● 我的代碼可以在任何地方運行嗎？
● 我可以同時使用 GPU 和 CPU 嗎？

所有這些問題的答案都很簡單。首先，我們必須考慮數據大小和算法結構。如果我們的數組很小或者我們的算法是高度順序的，那么 CPU 可能是最佳選擇。但是，一旦我們的數據增長，或者我們的計算可以被描述為“對每一行執(zhí)行相同的操作”，GPU 就會脫穎而出。如果我們堅持使用現代 Python 庫，除非我們愿意，否則我們不需要編寫低級 CUDA C 代碼。我們將獲得熟悉的、感覺像 NumPy 的高級 API。本書將向您展示：逐步講解如何從基本數據結構過渡到啟動我們自己的 GPU 內核。

我們還需要兼容的 GPU（通常選擇 NVIDIA 的 CUDA 庫）和合適的驅動程序。對于許多云提供商來說，只需點擊一下即可獲得 GPU 實例。即使是消費級筆記本電腦也通常配備支持 CUDA 的硬件。一旦我們的環(huán)境搭建完成，工作流程就會像任何 Python 項目一樣無縫銜接。我們必須記住，許多工作負載使用 CPU 進行編排、復雜邏輯和數據準備，而 GPU 則負責處理繁重的工作。我們通常會將數據復制到 GPU，運行內核，然后再將結果復制回來。這種混合工作流程在科學計算和機器學習中很常見。

1.1.7 讓我們開始吧！

我們將通過對一個簡單的操作進行 CPU 和 GPU 的基準測試，邁出實際編碼的第一步。您將看到整個工作流程：用 Python 創(chuàng)建數據，在 CPU 上處理數據，然后使用 CuPy 或 PyCUDA 在 GPU 上運行相同的邏輯。您將比較運行時間，驗證正確性，并學習如何解讀我們的結果。整本書的基調是通過這種實踐方法，我們將掌握所需的知識。每個概念都與一個特定的編碼任務相關聯(lián)，我們可以在自己的機器上運行和修改這些任務。隨著學習的深入，我們將學習了解哪些操作需要外包，以及如何改進內存移動和內核配置。
您正處于一條實踐之路的起點，這條道路將改變我們用 Python 實現高性能計算的方式。我們將學習如何識別數據并行的機會、分析代碼，并在 CPU 和 GPU 之間輕松切換，從而使我們能夠自信地應對現代數據挑戰(zhàn)。如果您希望提高數據管道的速度、創(chuàng)建機器學習服務，或者只是想更快地分析數據，GPU 編程可以將我們的性能提升到一個全新的水平。我們將在實際代碼中，觀察這些想法如何轉化為實際的加速效果，從而繼續(xù)前進。

1.2 流式多處理器和核心概念

您已經看到了 GPU 加速的潛力，但要充分利用它，我們需要了解硬件內部的運作方式。與專注于高時鐘速度和復雜指令邏輯的 CPU 不同，GPU 會組織其資源以實現并行吞吐量。每個現代 GPU 的核心都是一組流式多處理器 (SM：Streaming Multiprocessors)。每個 SM 都像一個強大的引擎，旨在并行執(zhí)行數千個線程，并采用旨在隱藏延遲和保持數據傳輸的結構。

我們將深入剖析其結構。當我們觀察 GPU 時，我們會發(fā)現數十個 SM 并排排列。每個 SM 包含一組簡單的 CUDA 核心（基本計算單元）、特殊功能單元、寄存器、共享內存和 Warp 調度器。可以將 SM 視為一個獨立的微型處理器，它管理著自己的線程池。SM 的數量因 GPU 型號而異——入門級顯卡可能只有幾個，而數據中心 GPU 則擁有 80 個或更多。

1.2.1 SM 如何工作？

當我們啟動一個內核時，實際上是在告訴 GPU：“這是一個巨大的線程網格；請在所有線程上執(zhí)行相同的代碼。” GPU 會將我們的線程網格拆分成多個塊，每個塊最多包含數百個線程。然后，每個塊會被分配給一個 SM。接下來，事情變得有趣起來：每個 SM 可以同時執(zhí)行多個塊，具體取決于資源可用性（寄存器、共享內存等）。

來源：https://developer.nvidia.com/blog/cuda-refresher-getting-started-with-cuda

在 SM 中，線程進一步分組為 Warp（線程束）——由 32 個線程組成的線程組，它們以同步方式運行。每個周期，SM 的 Warp 調度器都會選擇一個或多個活動的 Warp，并將它們發(fā)送到 CUDA 核心進行執(zhí)行。一個 Warp 中的所有線程同時執(zhí)行同一條指令，但操作不同的數據。這是 SIMD（單指令多數據）編程的基礎，也是 GPU 在處理大型數組或圖像時如此高效的原因。

您可以想象一個大教室，每一行（warp）都遵循相同的學習計劃，但每個學生都在自己的工作表上學習。SM 的 warp 調度程序可以同時管理多個這樣的“行”，并在它們之間快速切換，以確保每個人都保持忙碌，即使某些線程需要等待內存或數據。

1.2.2 什么是占用率（occupancy）？

在 GPU 編程中，您會經常聽到“占用率”這個術語。占用率衡量我們對 GPU 并行資源的利用效率，具體來說，每個 SM 上有多少個處于活動狀態(tài)的 warp 相對于最大可能數量。更高的占用率意味著更多的 warp 準備運行，這有助于 SM 隱藏內存延遲。如果一個 warp 因等待數據而停滯，調度程序只需切換到另一個 warp 即可。
您可以通過在每個塊中啟動更多線程、減少每個線程的寄存器和共享內存使用量，或優(yōu)化內核以避免瓶頸來提高占用率。然而，更高的占用率并不總是能保證更好的性能——有時，我們會找到一個平衡資源使用率和吞吐量的最佳點。

1.2.3 Warp 調度和性能模式

每個 SM 中的 Warp 調度器就像一個用于并行執(zhí)行的空中交通管制員。它會根據可用性、就緒情況和資源限制來選擇運行哪些 Warp。當一個 Warp 正在等待內存或同步點時，調度器會快速切換入另一個 Warp，從而保持硬件繁忙。這種快速的上下文切換非常輕量，因為所有線程及其數據都駐留在 SM 上。
在處理高延遲操作（例如從全局內存讀取）時，您就能體會到這種架構的真正威力。Warp 調度器會同時處理數十個 Warp、重疊計算和數據傳輸，因此 SM 永遠不會空閑。如果我們構建代碼以提供足夠的并行工作，調度器就會為您處理延遲隱藏。

1.2.4 運行簡單的微基準測試

我們將進行實踐操作。您可以使用小型 PyCUDA 或 CuPy 內核探索 SM 和 Warp 調度概念。假設我們想要測量 GPU 可以并行處理的線程數，以及執(zhí)行時間如何隨不同的塊大小而變化。這不僅僅是一個學術練習——它是調優(yōu)代碼以實現高性能的第一步。
首先，設置一個包含數百萬個元素的數組。然后，編寫一個內核，對每個元素執(zhí)行一個簡單的操作（例如，加一個常量）。啟動內核，并使用不同的塊大小（例如，每個塊 32、64、128、256、512 和 1024 個線程），并測量每個塊的執(zhí)行時間。我們可能會觀察到某些配置的運行速度比其他配置快得多，這取決于它們與我們 GPU 的 SM 架構的“匹配程度”。通過使用 Python 查詢（CuPy 或 PyCUDA 提供設備屬性）打印出 SM 的數量和每個 SM 的最大線程數，我們可以將啟動配置與硬件的最佳性能相匹配。嘗試用較少的線程運行相同的操作——你的 GPU 將無法充分利用。如果每個塊的線程過多，我們可能會耗盡寄存器或共享內存的減少，導致占用率下降并損害性能。
在我們進行這些微基準測試的過程中，一些高性能模式變得清晰起來。首先，我們希望在不過度消耗資源的情況下最大化每個 SM 的線程數。其次，我們傾向于啟動足夠多的塊，以便每個 SM 保持繁忙，尤其是在更大的 GPU 上。第三，我們注意到內存合并的重要性：構建數據訪問，以便 Warp 中的線程從連續(xù)的內存地址讀取數據，從而提高吞吐量。當我們將這些模式付諸實踐時，即使是用于向量加法或直方圖計算等任務的簡單內核，其速度也會比簡單的實現快幾個數量級。
如果我們理解了 SM、Warp、占用率和調度程序，我們就可以自信地嘗試啟動參數和內核資源使用情況。我們不僅僅是編寫代碼，而是根據硬件的節(jié)奏調整程序。我們編寫的每個內核都提供了一個實時觀察決策效果的機會。

1.3 線程塊、網格和索引

我相信您現在已經聽說過 CUDA。這項技術使我們能夠在自己的代碼中充分利用 NVIDIA GPU 的強大功能。CUDA（統(tǒng)一計算設備架構：Compute Unified Device Architecture）提供了在 GPU 上啟動數千個輕量級線程的框架。使用 Python，我們可以使用 PyCUDA 或 CuPy 等庫來編寫編譯和運行 CUDA 內核的代碼，同時又能保持我們慣用的 Python 環(huán)境。本書將使用 Linux 作為默認設置，因為它對 CUDA 開發(fā)提供了強大的支持。

（來源：https://developer.nvidia.com/blog/cuda-refresher-cuda-programming-model）

1.3.1 線程層次結構

在 CUDA 中，并行工作使用兩級層次結構來組織：網格和塊。每次啟動內核時，我們都會指定要運行的塊數以及每個塊中要放置的線程數。網格中的每個線程執(zhí)行相同的代碼，但處理數據的不同部分。將我們的數據想象成一個大型電子表格。我們的想法是為電子表格的每個區(qū)域（或塊）分配一個團隊，并為該區(qū)域中的每個單元格（或線程）分配一個工作線程。CUDA 的架構使我們能夠處理海量數據集：一維、二維甚至三維數組，方法是將每個數據元素映射到網格中其自己的線程。

1.3.2 映射一維數組

為此，我們需要一個支持 CUDA 的 Python 環(huán)境。如果這是我們第一次嘗試，請繼續(xù)在虛擬環(huán)境中安裝 PyCUDA 或 CuPy，并確認我們的 GPU 可見。如果我們想要檢查，只需導入我們的庫并打印設備屬性即可。例如，使用 CuPy：

>>> import cupy as cp
>>> print(cp.cuda.runtime.getDeviceProperties(0))
{'name': b'NVIDIA GeForce RTX 4090', 'totalGlobalMem': 25241190400, 'sharedMemPerBlock': 49152, 'regsPerBlock': 65536, 'warpSize': 32, 'maxThreadsPerBlock': 1024, 'maxThreadsDim': (1024, 1024, 64), 'maxGridSize': (2147483647, 65535, 65535), 'clockRate': 2520000, 'totalConstMem': 65536, 'major': 8, 'minor': 9, 'textureAlignment': 512, 'texturePitchAlignment': 32, 'multiProcessorCount': 128, 'kernelExecTimeoutEnabled': 0, 'integrated': 0, 'canMapHostMemory': 1, 'computeMode': 0, 'maxTexture1D': 131072, 'maxTexture2D': (131072, 65536), 'maxTexture3D': (16384, 16384, 16384), 'concurrentKernels': 1, 'ECCEnabled': 0, 'pciBusID': 172, 'pciDeviceID': 0, 'pciDomainID': 0, 'tccDriver': 0, 'memoryClockRate': 10501000, 'memoryBusWidth': 384, 'l2CacheSize': 75497472, 'maxThreadsPerMultiProcessor': 1536, 'isMultiGpuBoard': 0, 'cooperativeLaunch': 1, 'cooperativeMultiDeviceLaunch': 1, 'deviceOverlap': 1, 'maxTexture1DMipmap': 32768, 'maxTexture1DLinear': 268435456, 'maxTexture1DLayered': (32768, 2048), 'maxTexture2DMipmap': (32768, 32768), 'maxTexture2DLinear': (131072, 65000, 2097120), 'maxTexture2DLayered': (32768, 32768, 2048), 'maxTexture2DGather': (32768, 32768), 'maxTexture3DAlt': (8192, 8192, 32768), 'maxTextureCubemap': 32768, 'maxTextureCubemapLayered': (32768, 2046), 'maxSurface1D': 32768, 'maxSurface1DLayered': (32768, 2048), 'maxSurface2D': (131072, 65536), 'maxSurface2DLayered': (32768, 32768, 2048), 'maxSurface3D': (16384, 16384, 16384), 'maxSurfaceCubemap': 32768, 'maxSurfaceCubemapLayered': (32768, 2046), 'surfaceAlignment': 512, 'asyncEngineCount': 2, 'unifiedAddressing': 1, 'streamPrioritiesSupported': 1, 'globalL1CacheSupported': 1, 'localL1CacheSupported': 1, 'sharedMemPerMultiprocessor': 102400, 'regsPerMultiprocessor': 65536, 'managedMemory': 1, 'multiGpuBoardGroupID': 0, 'hostNativeAtomicSupported': 0, 'singleToDoublePrecisionPerfRatio': 64, 'pageableMemoryAccess': 1, 'concurrentManagedAccess': 1, 'computePreemptionSupported': 1, 'canUseHostPointerForRegisteredMem': 1, 'sharedMemPerBlockOptin': 101376, 'pageableMemoryAccessUsesHostPageTables': 0, 'directManagedMemAccessFromHost': 0, 'uuid': b'\xea\xce)\xa0\xbe\xf8\x9d\xfe\x06w\x96\xdb\x89\x84,\xc4', 'luid': b'', 'luidDeviceNodeMask': 0, 'persistingL2CacheMaxSize': 51904512, 'maxBlocksPerMultiProcessor': 24, 'accessPolicyMaxWindowSize': 134213632, 'reservedSharedMemPerBlock': 1024}

之后，我們將從一個簡單的一維數組開始。假設我們有一個包含一百萬個元素的數組，并且我們希望每個線程對其分配的元素進行平方。啟動內核時，我們會確定每個塊的線程數（例如 256），以及所需的塊數（num_elements //threads_per_block + 1）。
在 CUDA 內核中，每個線程都會計算出其唯一的索引：

int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < num_elements) {
   output[idx] = input[idx] * input[idx];
}

blockIdx.x * blockDim.x + threadIdx.x 的計算為每個線程在整個網格中提供了唯一的全局索引。 if 語句是我們的邊界檢查——某些線程的索引可能超出了數組的末尾，因此我們阻止它們越界寫入。
在 Python 端，運行內核后，我們將輸出數組從設備復制回主機并檢查結果：

import cupy as cp
num_elements = 1_000_000
input_array = cp.arange(num_elements, dtype=cp.float32)
output_array = cp.empty_like(input_array)
threads_per_block = 256
blocks_per_grid = (num_elements + threads_per_block - 1) // threads_per_block
kernel_code = '''
extern "C" __global__
void square(const float *input, float *output, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
   if (idx < n) {
       output[idx] = input[idx] * input[idx];
   }
}
'''
module = cp.RawModule(code=kernel_code)
square_kernel = module.get_function('square')
square_kernel((blocks_per_grid,), (threads_per_block,),
             (input_array, output_array, num_elements))
# Copy results to host and verify
result = cp.asnumpy(output_array)
assert (result == (cp.asnumpy(input_array) ** 2)).all()
print("1D kernel executed successfully!")

我們剛剛將一維數組映射到 CUDA 的網格和塊系統(tǒng)上。

1.3.3 映射二維數組

現在，假設我們要處理一個二維數組，例如圖像或矩陣。您需要將行和列映射到兩個網格維度。 CUDA 支持多維塊和網格，允許每個線程處理單個像素或矩陣條目。
現在，在我們的內核中，我們需要計算行和列的索引：

int row = blockIdx.y * blockDim.y + threadIdx.y;
int col = blockIdx.x * blockDim.x + threadIdx.x;
if (row < height && col < width) {
    output[row * width + col] = input[row * width + col] + 1.0f;
}

在 Python 中，我們可以調整啟動配置：

import cupy as cp

height, width = 1024, 1024
input_array = cp.random.rand(height, width).astype(cp.float32)
output_array = cp.empty_like(input_array)
block = (16, 16)
grid = ((width + block[0] - 1) // block[0],
       (height + block[1] - 1) // block[1])
kernel_code = '''
extern "C" __global__
void increment(const float *input, float *output, int width, int height) {
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    if (row < height && col < width) {
       output[row * width + col] = input[row * width + col] + 1.0f;
   }
}
'''
module = cp.RawModule(code=kernel_code)
increment_kernel = module.get_function('increment')
increment_kernel(grid, block, (input_array, output_array, width, height))
result = cp.asnumpy(output_array)
assert (result == cp.asnumpy(input_array) + 1).all()
print("2D kernel executed and verified successfully!")

之后，每個線程計算其行和列。邊界檢查確保如果網格大小與數組尺寸不完全匹配，我們不會越過邊界。

您可以將此模式擴展到三維數據，例如體積圖像或模擬網格。 CUDA 允許我們指定三維塊和網格，每個線程計算 (z, y, x) 索引來處理其自身的體素或數據點。關鍵始終如一：將邏輯數據位置（無論是一維、二維還是三維）映射到 CUDA 的線程和塊索引，檢查邊界，并確保每個線程僅在有效數據范圍內工作。

了解了這一點后，我們就可以開始對任何數據形狀使用自定義內核了。將 Python 數組映射到 CUDA 網格和塊、處理邊界以及驗證輸出的基礎知識已經包含在我們的工具包中。這些模式幾乎會體現在我們構建的每個 GPU 應用程序中，為實際的高性能 Python 代碼奠定基礎。

1.4 主機與設備交互概述

您將親身體驗我們的 Python 代碼如何與 GPU 通信。此過程包含幾個核心步驟：內存分配、數據傳輸、啟動內核以及了解所涉及的不同內存空間。您已經安裝了 CuPy，因此可以開始將理論付諸實踐。我們將通過編寫一個涵蓋所有基本知識的簡潔而完整的腳本來構建我們的直覺。

1.4.1 GPU 上的內存分配

在 CuPy 中，當我們分配數組時，實際上是直接在 GPU 上預留內存，而不是在計算機的主內存上。這是一個巨大的思維轉變。如果我們使用 CuPy，我們創(chuàng)建的數組將保存在設備上。
這里有一個簡單的方法：

import cupy as cp
# Allocate an array of 10,000 floats on the GPU
gpu_array = cp.zeros(10_000, dtype=cp.float32)
print("Array allocated on GPU:", gpu_array)

這一行代碼為我們分配了一塊 GPU 內存。與 NumPy 數組不同，這些數組不能直接用于僅使用 CPU 的函數。

1.4.2 主機到設備，以及從設備到主機

主要的 Python 腳本在 CPU（“主機”）上運行，但繁重的工作在 GPU（“設備”）上進行。要使用 GPU 加速，我們首先需要將數據從主機移動到設備。
因此，我們首先創(chuàng)建一個 NumPy 數組，然后將其傳輸到 GPU。

import numpy as np
import cupy as cp

# Create data on the host
host_data = np.arange(10_000, dtype=np.float32)
# Transfer data to the GPU (device)
device_data = cp.asarray(host_data)
print("Data transferred to GPU.")

# Bring the data back from device to host
result_on_host = cp.asnumpy(device_data)
print("Data transferred back to host. First five elements:", result_on_host[:5])

您經常會發(fā)現自己需要來回移動數據，尤其是在我們需要驗證輸出、保存結果或與其他需要 NumPy 數組的庫交互時。

1.4.3 啟動簡單內核

CuPy 的一大亮點在于它允許我們直接在 Python 中編寫自定義 CUDA 內核。假設有一個數組，我們想將每個元素乘以 2。你可以使用 CuPy 的 RawKernel 接口來編寫這個內核。

import cupy as cp

# 1. 內核代碼定義
kernel_code = r'''
extern "C" __global__
void multiply_by_two(float* data, int n) {
    int idx = blockDim.x * blockIdx.x + threadIdx.x;
    if (idx < n) {
        data[idx] *= 2.0f;
    }
}
'''

# 2. 編譯內核
module = cp.RawModule(code=kernel_code)
multiply_by_two = module.get_function('multiply_by_two')

# 3. 設置參數和數據
n = 10_000
threads_per_block = 256
blocks_per_grid = (n + threads_per_block - 1) // threads_per_block

# 4. 創(chuàng)建并初始化GPU數組
gpu_array = cp.random.rand(n, dtype=cp.float32)

print("Original array (first 5 elements):", gpu_array[:5].get()) # .get() 將數據從GPU傳回CPU

# 5. 啟動內核
multiply_by_two((blocks_per_grid,), (threads_per_block,), (gpu_array, n))

# 6. 等待內核執(zhí)行完畢并驗證結果
cp.cuda.stream.get_current_stream().synchronize()
print("Kernel launched and synchronized.")
print("Modified array (first 5 elements):", gpu_array[:5].get()) # 再次將數據傳回CPU，查看結果

1.4.4 理解 GPU 內存空間

所以到目前為止，您已經使用了兩個主要的“內存空間”：
● 主機內存（您的主系統(tǒng) RAM），NumPy 數組存放于此，Python 腳本也從這里開始。
● 另一個是設備內存（GPU），CuPy 數組在此分配，內核也在此讀寫數據。
GPU 編程的魅力在于了解何時移動數據以及每個操作在何處進行。CuPy 數組上的大多數計算都完全在設備上進行，從而保持快速運行并避免不必要的傳輸。
此外，高級內核還可以訪問：
● 共享內存，這是一個在塊中的線程之間共享的小型、快速的內存區(qū)域，對于需要協(xié)作的算法非常有用。
● 寄存器，每個線程私有的超高速存儲空間。
● 常量內存和紋理內存，它們是用于只讀或空間相干數據的專用空間，我們將在后面的章節(jié)中探討。
有了這些，我們現在有了一個完整的工作流程，如下所示：
● 在 GPU 上分配數組
● 在主機和設備之間移動數據
● 編寫并啟動簡單的內核
● 觀察和管理 GPU 內存使用情況
每個支持 GPU 加速的項目都會遵循這些步驟。隨著我們水平的提升，我們希望盡可能多地在設備上進行計算，并且僅將主機-設備傳輸用于輸入和最終輸出。掌握這些知識后，您就可以將我們的 GPU 工作流程提升到一個新的水平。

參考資料

• 軟件測試精品書籍文檔下載持續(xù)更新 https://github.com/china-testing/python-testing-examples 請點贊，謝謝！
• 本文涉及的python測試開發(fā)庫 謝謝點贊！ https://github.com/china-testing/python_cn_resouce
• python精品書籍下載 https://github.com/china-testing/python_cn_resouce/blob/main/python_good_books.md
• Linux精品書籍下載 http://www.rzrgm.cn/testing-/p/17438558.html
• python八字排盤 https://github.com/china-testing/bazi
• 聯(lián)系方式：釘ding或V信: pythontesting

1.5 簡單向量加法內核

我們即將構建第一個自定義 GPU 內核，用于執(zhí)行向量加法等基本任務。每個新加入團隊的 GPU 程序員都必須完成這項練習才能加入團隊。我們將學習如何使用 CuPy 和 PyCUDA 創(chuàng)建、啟動和驗證一個用于將兩個向量相加的內核。我們還會將 GPU 上的結果與 CPU 上的結果進行比較，以確保一切正常。

1.5.1 安裝 PyCUDA

我們已經準備好了 CuPy，它非常適合大多數高級 GPU 陣列工作。如果您想要更多控制權并直接訪問 CUDA C 語言，可以使用 PyCUDA。它允許我們編寫原始 CUDA 內核，并賦予我們對內存管理、內核啟動和設備上下文的強大控制權，所有這些都可以通過 Python 實現。
首先，讓我們在 Python 環(huán)境中安裝 PyCUDA。

pip install pycuda
``

安裝完成后，我們可以導入 PyCUDA 模塊并立即開始使用其功能。

### 1.5.2 主機和設備上的向量

- CuPy

我們首先使用 NumPy（CPU）和 CuPy（GPU）設置輸入數據：
```python
import numpy as np
import cupy as cp

# 設置向量大小
N = 1_000_000
# 在主機（CPU）上創(chuàng)建兩個隨機輸入向量
a_host = np.random.rand(N).astype(np.float32)
b_host = np.random.rand(N).astype(np.float32)
# 計算CPU上的結果，用于后續(xù)驗證
c_cpu = a_host + b_host
# 使用 CuPy 將輸入向量傳輸到 GPU
a_gpu = cp.asarray(a_host)
b_gpu = cp.asarray(b_host)

# CUDA C++ 內核代碼
kernel_code = r'''
// 聲明為全局函數，可以在GPU上執(zhí)行
extern "C" __global__
// 定義向量相加函數，接受三個數組指針和數組大小
void vector_add(const float* a, const float* b, float* c, int n) {
    // 計算當前線程在整個網格中的全局唯一索引
    int idx = blockDim.x * blockIdx.x + threadIdx.x;
    // 邊界檢查，確保線程索引不超過數組大小
    if (idx < n) {
        // 執(zhí)行向量加法操作
        c[idx] = a[idx] + b[idx];
    }
}
'''

# 在GPU上創(chuàng)建用于存放結果的空數組
c_gpu = cp.empty_like(a_gpu)
# 定義每個塊的線程數
threads_per_block = 256
# 計算需要的塊數，確保所有元素都能被處理到
blocks_per_grid = (N + threads_per_block - 1) // threads_per_block
# 編譯CUDA內核代碼
module = cp.RawModule(code=kernel_code)
# 從編譯模塊中獲取內核函數
vector_add = module.get_function('vector_add')
# 啟動內核，并傳入網格、塊的維度以及函數參數
vector_add((blocks_per_grid,), (threads_per_block,), (a_gpu, b_gpu, c_gpu, N))

# 將GPU上的結果數組傳輸回主機（CPU）
c_result = cp.asnumpy(c_gpu)
# 比較GPU和CPU的結果，檢查它們是否幾乎相等
print("Are GPU and CPU results equal?", np.allclose(c_result, c_cpu))

• 使用 CuPy 運行內核

import numpy as np
import pycuda.driver as cuda
import pycuda.autoinit
from pycuda.compiler import SourceModule

# 設置向量大小
N = 1_000_000
# 在主機（CPU）上創(chuàng)建兩個隨機輸入向量
a_host = np.random.rand(N).astype(np.float32)
b_host = np.random.rand(N).astype(np.float32)

# 計算CPU上的結果，用于后續(xù)驗證
c_cpu = a_host + b_host

# 在GPU上分配內存
# pycuda.mem_alloc 分配的內存需要手動管理
a_gpu = cuda.mem_alloc(a_host.nbytes)
b_gpu = cuda.mem_alloc(b_host.nbytes)
c_gpu = cuda.mem_alloc(a_host.nbytes)

# 將數據從主機（CPU）傳輸到設備（GPU）
cuda.memcpy_htod(a_gpu, a_host)
cuda.memcpy_htod(b_gpu, b_host)

# CUDA C++ 內核代碼
kernel_code = """
__global__ void vector_add(const float* a, const float* b, float* c, int n) {
    int idx = blockDim.x * blockIdx.x + threadIdx.x;
    if (idx < n) {
        c[idx] = a[idx] + b[idx];
    }
}
"""

# 編譯內核代碼
# pycuda.compiler.SourceModule 負責編譯
mod = SourceModule(kernel_code)
vector_add = mod.get_function("vector_add")

# 定義每個塊的線程數和需要的塊數
threads_per_block = 256
blocks_per_grid = (N + threads_per_block - 1) // threads_per_block

# 啟動內核
# 注意參數的傳遞方式：(grid_dim, block_dim, shared_mem, stream, *args)
# 這里我們沒有使用共享內存和流，所以為0和None
vector_add(a_gpu, b_gpu, c_gpu, np.int32(N),
           grid=(blocks_per_grid, 1, 1),
           block=(threads_per_block, 1, 1))

# 從設備（GPU）將結果傳輸回主機（CPU）
c_result = np.empty_like(a_host)
cuda.memcpy_dtoh(c_result, c_gpu)

# 比較GPU和CPU的結果，檢查它們是否幾乎相等
print("Are GPU and CPU results equal?", np.allclose(c_result, c_cpu))

# 釋放GPU內存
a_gpu.free()
b_gpu.free()
c_gpu.free()

另外一種更簡潔的實現：

import numpy as np
import pycuda.driver as cuda
import pycuda.autoinit
import pycuda.gpuarray as gpuarray
from pycuda.compiler import SourceModule

# --- 1. 定義缺失的變量和代碼 ---
kernel_code = """
__global__ void vector_add(const float* a, const float* b, float* c, int n) {
    int idx = blockDim.x * blockIdx.x + threadIdx.x;
    if (idx < n) {
        c[idx] = a[idx] + b[idx];
    }
}
"""

# 設置向量大小
N = 1_000_000
# 在主機（CPU）上創(chuàng)建兩個隨機輸入向量
a_host = np.random.rand(N).astype(np.float32)
b_host = np.random.rand(N).astype(np.float32)

# 計算CPU結果用于驗證
c_cpu = a_host + b_host

threads_per_block = 256
blocks_per_grid = (N + threads_per_block - 1) // threads_per_block
# ------------------------------------

# Transfer input data to GPU
a_gpu_py = gpuarray.to_gpu(a_host)
b_gpu_py = gpuarray.to_gpu(b_host)
c_gpu_py = gpuarray.empty_like(a_gpu_py)

mod = SourceModule(kernel_code)
vector_add_func = mod.get_function("vector_add")

# --- 2. 修正內核啟動參數 ---
vector_add_func(
    a_gpu_py, b_gpu_py, c_gpu_py, np.int32(N),
    block=(threads_per_block, 1, 1), grid=(blocks_per_grid, 1, 1)
)

# 獲取GPU結果
c_result_py = c_gpu_py.get()
print("PyCUDA GPU result matches CPU?", np.allclose(c_result_py, c_cpu))

現在，我們應該再次看到 True。兩個庫都給出了相同的結果，證實我們的內核可以正常工作！這樣，我們就可以在 CuPy 和 PyCUDA 中編寫、編譯并啟動一個自定義 CUDA 內核了。我們以閃電般的速度執(zhí)行了元素向量加法，并驗證了結果與我們的 CPU 計算結果一致。這是所有后續(xù) GPU 編程的基礎，因為每個復雜的操作都是基于這些模式構建的。

1.6 小結

本章講解了所有必要的基礎模塊，為實際的 GPU 編程奠定了堅實的基礎。我們首先了解了現實世界對高性能數據處理的需求，了解了傳統(tǒng) CPU 在現代工作負載下經常遇到的瓶頸，并理解了擁有數千個輕量級核心的 GPU 為何能夠加速機器學習推理和大規(guī)模分析等任務。我們研究了流多處理器的設計，了解了 SM、warp 和占用率指標如何讓您高效地并行運行。我們還學習了 GPU 隱藏內存延遲并保持所有任務繁忙的技巧。

然后，我們深入研究了將多維數組映射到 CUDA 網格和塊。我們掌握了線程層次結構、索引計算和穩(wěn)健邊界處理的概念。我們用 Python 完成了整個 GPU 數據處理工作流程，從使用 CuPy 設置設備內存，到在主機和設備之間移動數據、啟動自定義內核，再到將結果復制回來以驗證所有內容。我們嘗試使用 CuPy 和 PyCUDA 編寫并運行了我們的第一個自定義 CUDA C 內核，并經歷了從在 CPU 上準備數據到驗證 GPU 輸出的整個過程。憑借這些技能和習慣，我們將在本書的其余部分深入探索 GPU 編程的更高級、更具創(chuàng)造性的用法。