CUDA性能優化實戰：7個步驟讓并行歸約算法提升10倍效率

本文深入探討了一個經典的并行計算算法——并行歸約（Parallel Reduction）的性能優化過程，通過七個漸進式的優化步驟，展示了如何將算法性能提升至極致。這項研究基于Mark Harris在NVIDIA網絡研討會中提出的優化方法，在重現這些優化技術的同時，進一步簡化了概念闡述以便于理解。配套的GitHub代碼庫提供了完整的實現細節，為讀者深入研究提供了詳實的技術支撐。

算法原理分析

并行歸約算法是CUDA編程中的一個重要數據并行原語，其核心思想是利用GPU的線程層次結構對向量、矩陣或張量進行并行計算。該算法通過sum()、min()、max()或avg()等操作對數據進行聚合處理。本文將重點使用sum()操作進行數據集歸約。盡管這些操作在概念上相對簡單，但它們在眾多應用場景中發揮著關鍵作用，因此需要高度優化以避免成為性能瓶頸。

在并行化實現過程中，算法采用基于樹的方法，計算任務分布在GPU的各個線程塊中。這里面臨一個核心技術挑戰：如何在線程塊之間高效傳遞部分計算結果？最直觀的解決方案是采用全局同步機制——讓各個塊完成計算后進行全局同步，然后遞歸繼續處理。CUDA架構并不支持全局同步，主要原因是硬件成本過高，且會限制程序員只能使用少量線程塊以避免死鎖，從而顯著降低整體計算效率。

基于樹的歸約 | 來源：NVIDIA

解決線程塊間部分結果通信問題的實用方法是采用內核分解技術。內核分解將大規模的內核任務分解為多個較小的、可管理的子任務，這些子任務可以在不同的線程或塊中獨立執行。這種方法最大限度地減少了硬件和軟件開銷，實現了更靈活高效的GPU資源利用，同時降低了同步需求并提升了整體計算性能。

內核分解 | 來源：NVIDIA

性能評估指標體系

算法性能評估基于兩個關鍵維度：執行時間和帶寬利用率。這些指標能夠準確反映GPU資源利用程度，本質上衡量系統是否達到了峰值性能。我們的優化目標是實現GPU峰值性能，通過計算性能（GFLOP/s）和內存性能（GB/s）兩個方面的指標進行量化評估。

為了實現這些指標的優化，需要重點關注數據訪問模式和計算瓶頸識別兩個核心方面。具體而言，需要評估如何提升數據讀寫效率，以及如何使計算過程更加快速和高效。在GPU編程中，理想的計算實現不僅要追求高速度，更要確保大部分線程都能有效參與工作。

REDUCE-0：交替尋址基礎實現

作為優化的起點，首先實現最基礎的并行歸約方法。這種樸素的并行化方法需要確定訪問存儲元素的地址空間模式，檢索相應元素，通過求和操作組合這些元素，并在不同線程上遞歸重復此過程以實現操作的并行化。

交替尋址技術的核心是訪問和組合位于當前線程處理段中間位置的地址空間。以包含1024個整數的數組為例，如果每個塊使用256個線程，每個線程從不同起點開始，每次跳過256個元素進行處理。線程0將依次處理元素0、256、512和768，每次將當前元素與位于其負責數組段中間位置的另一個元素進行組合。因此，線程0會將元素0與元素128組合，元素256與384組合，元素512與640組合，元素768與896組合。這個過程將遞歸進行直到獲得最終結果。

這種方法在簡化線程間同步的同時，確保所有線程都能積極參與并行數據歸約，從而實現更加均衡的負載分配和高效的歸約計算。

交替尋址 | 來源：NVIDIA

 // 歸約 0 – 交替尋址  
__global__ void reduce0(int *g_in_data, int *g_out_data){  extern __shared__ int sdata[];  // 存儲在共享內存中  // 每個線程從全局內存加載一個元素到共享內存  unsigned int tid = threadIdx.x;  unsigned int i = blockIdx.x * blockDim.x + threadIdx.x;  sdata[tid] = g_in_data[i];  __syncthreads();  // 歸約方法 -- 在共享內存中進行，因為sdata存儲在那里  for(unsigned int s = 1; s < blockDim.x; s *= 2){  if (tid % (2 * s) == 0) {  sdata[tid] += sdata[tid + s];     }  __syncthreads();  }  if (tid == 0){  g_out_data[blockIdx.x] = sdata[0];  }  }

實現過程包含六個關鍵步驟。首先，根據線程ID和塊大小為每個線程分配起始索引。接著，每個線程將對應元素從全局內存加載到共享內存。然后同步塊內所有線程以確保數據加載完成。在共享內存中執行歸約操作時，每個線程將其值與另一個線程在計算偏移量處的值相加，該偏移量在每個后續步驟中減半。每個歸約步驟后需要再次同步線程以保證數據完整性。最后，每個塊中的第一個線程將歸約結果寫入全局內存的輸出數組。

結果

REDUCE-0 結果

性能瓶頸分析

雖然這種方法為并行編程奠定了良好基礎，但仍存在明顯的性能瓶頸。從計算和內存兩個維度分析，可以識別出以下效率問題。

在計算方面，主要瓶頸來源于模運算符（%）的使用。該運算符在計算上開銷很大，因為它涉及除法操作——這是底層硬件上最慢的操作之一。在內核中頻繁執行該操作會嚴重影響性能。此外交替尋址模式導致warp高度發散，因為同一warp內的線程由于條件判斷需要執行不同的代碼路徑。這種路徑發散導致warp停滯，等待其他線程完成，嚴重降低了執行效率。

在內存方面，由于warp發散導致的內存訪問模式次優。每個線程訪問分布在整個數組中的數據元素，使得內存訪問呈現分散且非連續的特點，導致內存帶寬利用效率低下和數據檢索延遲增高。這種分散的訪問模式可能產生多個緩慢的內存事務，而非單個高效的事務，因此無法充分利用GPU的內存帶寬能力。

REDUCE-1：改進的交替尋址

針對第一版實現中的計算效率問題，這個版本對尋址方式進行了優化。雖然基本的尋址邏輯保持不變，但在構建歸約函數時消除了模運算符和發散條件的使用。通過重構索引計算方式（

int index = 2 * s * tid;

），REDUCE-1確保每個線程能夠一致地執行操作，無需檢查相對于步長的位置，從而消除了warp內的發散現象。

這種調整使得warp中的所有線程都遵循相同的執行路徑，顯著提升了warp執行效率。移除模運算符進一步提升了性能，避免了GPU上運行緩慢的除法相關運算。

 // 歸約 1 – 無分支發散和%運算的交替尋址  
__global__ void reduce1(int *g_in_data, int *g_out_data){  extern __shared__ int sdata[];  // 存儲在共享內存中  // 每個線程從全局內存加載一個元素到共享內存  unsigned int tid = threadIdx.x;  unsigned int i = blockIdx.x * blockDim.x + threadIdx.x;  sdata[tid] = g_in_data[i];  __syncthreads();  // 歸約方法 -- 在共享內存中進行  for(unsigned int s = 1; s < blockDim.x; s *= 2){  // 注意步長為 s *= 2：這導致交替尋址  int index = 2 * s * tid;    // 現在我們不需要if條件的發散分支  if (index + s < blockDim.x)  {  sdata[index] += sdata[index + s];   // s用來表示將要合并的偏移量  }  __syncthreads();  }  if (tid == 0){  g_out_data[blockIdx.x] = sdata[0];  }  }

結果

REDUCE-1 結果

新問題的出現

盡管REDUCE-1在計算效率和執行一致性方面相比REDUCE-0有顯著改進，但引入了一個新的性能問題：共享內存庫沖突（Bank Conflicts）。當多個線程同時嘗試訪問同一內存庫的數據時，就會發生這種沖突，導致原本可以并行執行的內存訪問被強制串行化。

從REDUCE-0到REDUCE-1的轉換過程中，雖然提升了算法的計算效率，但并未解決內存相關問題，反而通過引入步長機制創造了更多內存訪問沖突。步長方法使得線程嘗試訪問相同的共享內存地址。REDUCE-0將線程分散在充當邊界的間隔內，將線程訪問限制在這些邊界內，從而減少了沖突機會。而REDUCE-1依賴步長并移除了這些邊界，導致庫沖突和進程串行化。

由于每個內存庫每個周期只能處理一次訪問，當多個訪問指向同一庫時必須進行串行化處理，這有效降低了內存操作的吞吐量。這種串行化抵消了通過消除warp發散獲得的部分性能提升，在較大線程塊中可能成為顯著的性能瓶頸。

REDUCE-2：順序尋址優化

為了解決內存訪問沖突問題，這個版本采用了更高效的尋址技術。與讓線程訪問間隔較遠元素的交替尋址不同，順序尋址讓每個線程處理連續的數據元素。

在1024元素、每塊256線程的示例中，線程0將訪問連續的元素0、1、2、3，而非間隔較遠的0、256、512、768。線程0依次組合元素0和1、然后處理元素2，以此類推進行遞歸處理。這種方法充分利用了空間局部性原理，通過提升緩存效率來避免庫沖突。該算法具有線性特征，最大限度地減少了會增加等待時間的同步需求。

順序尋址 | 來源：NVIDIA

這種變化通過將內存訪問模式與GPU對連續內存訪問的偏好更緊密對齊，顯著改善了內存訪問效率。通過訪問相鄰的內存位置，REDUCE-2降低了緩存未命中和內存庫沖突的概率，提升了內存帶寬利用效率，并改善了歸約操作的整體性能。

 // 歸約 2 – 順序尋址  
__global__ void reduce2(int *g_in_data, int *g_out_data){  extern __shared__ int sdata[];  // 存儲在共享內存中  // 每個線程從全局內存加載一個元素到共享內存  unsigned int tid = threadIdx.x;  unsigned int i = blockIdx.x * blockDim.x + threadIdx.x;  sdata[tid] = g_in_data[i];  __syncthreads();  // 歸約方法 -- 在共享內存中進行  for(unsigned int s = blockDim.x/2; s > 0; s >>= 1){  // REDUCE2 -- 查看上面的反向循環  if (tid < s){   // 然后，我們檢查線程ID來進行計算  sdata[tid] += sdata[tid + s];  }  __syncthreads();  }  if (tid == 0){  g_out_data[blockIdx.x] = sdata[0];  }  }

該方法的主要技術創新包括用反向循環結構結合基于線程ID的索引替換了步長索引機制，從根本上改變了歸約過程中的數據處理方式。反向循環從最高可能步長

s = blockDim.x / 2

開始歸約，每次迭代將步長減半。這意味著線程首先處理待求和數據間的最大間隙，快速減少需要處理的數據總量。基于線程ID的索引使每個線程使用其ID來訪問連續的數據點對而非分散的數據點，簡化了訪問模式并最小化了內存延遲。隨著步長的減小，線程組合相鄰元素，優化了內存使用并提升了數據吞吐量。

結果

REDUCE-2 結果

線程利用率問題

這種方法基本上解決了內存沖突問題。在解決了明顯的計算和內存問題后，需要進一步提升算法的智能化程度以獲得更好的性能表現。

當前面臨的主要問題是在第一個循環迭代中有一半的線程處于空閑狀態，這造成了資源浪費并未能充分利用GPU的計算能力。在1024元素的示例中，循環第一次迭代時

s=blockDim.x/2

（即

s=512

），條件

if (tid < s)

將活躍計算限制在塊的前512個線程。這意味著雖然這512個線程在積極地對元素對求和（例如

sdata[tid]

與

sdata[tid + 512]

），剩余的512個線程卻處于空閑狀態，對計算沒有任何貢獻。這種在每個后續迭代中將活躍線程數量減半的模式持續到歸約完成，從512減少到256，然后是128、64、32等。這種快速的線程數量衰減導致GPU能力的顯著浪費，特別是在初始迭代中只有一小部分可用線程參與工作。

解決方案是在數據加載到共享內存的同時進行第一次計算操作。

REDUCE-3：加載時預歸約

為了充分利用空閑線程并提升計算效率，在從全局內存向共享內存加載元素的同時執行第一次計算操作。這種方法能夠在加載過程中將兩個元素歸約為一個，從而將需要處理的數據塊數量減半。

具體實現中，在1024元素、256線程的配置下，每個線程將前兩個元素的和加載到共享內存中。線程0處理元素0和1，線程1處理元素2和3，以此類推。這樣可以將數據塊數量和共享內存長度都減半到512。代碼的其余部分與REDUCE-2完全相同，這意味著第一次迭代仍然會激活512個線程開始歸約操作，因為

s=blockDim.x/2 = 512

。這種方法讓更多線程參與有效工作，避免了計算資源的浪費。

 // 歸約 3 – 加載時首次加法  
__global__ void reduce3(int *g_in_data, int *g_out_data){  extern __shared__ int sdata[];  // 存儲在共享內存中  // 每個線程從全局內存加載一個元素到共享內存  unsigned int tid = threadIdx.x;  unsigned int i = blockIdx.x*(blockDim.x*2) + threadIdx.x;  sdata[tid] = g_in_data[i] + g_in_data[i+blockDim.x];  __syncthreads();  // 歸約方法 -- 在共享內存中進行  for(unsigned int s = blockDim.x/2; s > 0; s >>= 1){  // 查看上面的反向循環  if (tid < s){   // 然后，我們檢查tid來進行計算  sdata[tid] += sdata[tid + s];  }  __syncthreads();  }  if (tid == 0){  g_out_data[blockIdx.x] = sdata[0];  }  }

實現過程中包含三個關鍵技術變更。首先，在從全局內存加載元素到共享內存時同時進行初始歸約步驟：

sdata[tid] = g_in_data[i] + g_in_data[i+blockDim.x]

。其次，修改索引

的計算方式為

unsigned int i = blockId.x * (blockDim.x*2) + threadId.x

，因為每個線程現在同時處理兩個輸入，需要將每個塊覆蓋的有效索引范圍擴大一倍。最后，在主函數中修改內核調用方式，將執行配置設置為

int num_blocks = (n + (2*blockSize) - 1 / (2*blockSize)

，這樣可以將分配給內核的塊數量減半，同時保持代碼的正確性。

結果

REDUCE-3 結果

指令開銷瓶頸識別

當前方法表現良好，但仍有進一步優化的空間。通過分析性能指標發現，在Tesla T4上約41 GB/s的帶寬使用率表明我們并未達到或耗盡帶寬上限。另一方面，歸約操作具有低算術強度的特征，意味著我們也不受計算能力限制。

由于既不受帶寬限制也不受計算限制，還存在第三個潛在瓶頸：指令開銷。這包括GPU執行的所有輔助指令，這些指令不直接參與數據加載、存儲或歸約的主要算術操作。具體包括地址算術運算（計算下一個要加載的地址空間）和循環開銷（處理循環邏輯、條件判斷和迭代控制）。

針對這種瓶頸的優化策略是循環展開技術。

REDUCE-4：Warp級循環展開

首先分析REDUCE-3中的執行模式以理解優化的必要性。在1024元素的示例中，經過初始的元素對加載和相加后，256個線程處理512個元素。此時歸約過程中每個線程處理單個元素，線程活躍度呈現遞減模式：當

s = 256

時有256個活躍線程，當

s = 128

時有128個活躍線程，當

s = 64

時有64個活躍線程。

關鍵的優化點出現在

s = 32

時，此時有32個活躍線程。由于指令在warp內以SIMD方式同步執行，這帶來兩個重要特性：首先，無需使用

__syncthreads()

，因為所有線程在同一個warp中以鎖步方式工作；其次，無需

if (tid < s)

條件判斷，因為每個線程都需要執行相同的操作。因此可以安全地從這部分代碼中移除所有同步命令，顯著提升最終歸約階段的執行速度。

 // 添加這個函數來幫助展開  
__device__ void warpReduce(volatile int* sdata, int tid){  // 目標是讓所有warp避免無用的工作   sdata[tid] += sdata[tid + 32];  sdata[tid] += sdata[tid + 16];  sdata[tid] += sdata[tid + 8];  sdata[tid] += sdata[tid + 4];  sdata[tid] += sdata[tid + 2];  sdata[tid] += sdata[tid + 1];  
}  // 歸約 4 – 展開最后的Warp  
__global__ void reduce4(int *g_in_data, int *g_out_data){  extern __shared__ int sdata[];  // 存儲在共享內存中  // 每個線程從全局內存加載一個元素到共享內存  unsigned int tid = threadIdx.x;  unsigned int i = blockIdx.x*(blockDim.x*2) + threadIdx.x;  sdata[tid] = g_in_data[i] + g_in_data[i+blockDim.x];  __syncthreads();  // 只是將結束限制改為在s = 32之前停止  for(unsigned int s = blockDim.x/2; s > 32; s >>= 1){    // 查看上面的反向循環  if (tid < s){   // 然后，我們檢查tid來進行計算  sdata[tid] += sdata[tid + s];  }  __syncthreads();  }  // 添加這個以在s = 32時使用warpReduce  if (tid < 32){  warpReduce(sdata, tid);  }  if (tid == 0){  g_out_data[blockIdx.x] = sdata[0];  }  }

實現方法相對簡單：在

s = 32

之前停止主循環，并調用專門的

warpReduce

函數。該函數包含手寫的6次迭代，僅在設備端執行。同時需要使用

volatile

關鍵字確保實現的正確性。

結果

REDUCE-4 結果

擴展循環展開策略

這種優化取得了顯著的性能提升效果。既然循環展開如此有效，為什么不將這種技術擴展到更多的循環中？

REDUCE-5：完全循環展開

為了進一步擴展展開技術，需要在編譯時確定循環的總迭代次數。幸運的是，GPU將線程塊大小限制為512個線程，且通常使用2的冪次方配置。因此可以針對固定的塊大小進行完全展開，同時保持通用性。CUDA提供的C++模板參數支持使這種優化成為可能。

C++模板技術允許定義具有占位符的函數或類，這些占位符在編譯時被具體類型替換。通過使用模板參數來處理

blockSize

的變化，可以應對不同的展開需求。根據塊大小的不同，準備相應的switch case來處理特定的展開要求。完全展開技術消除了大部分歸約階段中不必要的循環和條件判斷，最小化了計算開銷。

通過編譯針對特定塊大小（如512、256和128）定制的內核版本，為每個變體優化其特定場景，剝離不必要的操作，最大化內存和計算資源效率。在具體實現中，主函數中將

blockSize

設置為256以簡化方法，同時包含了512、256和128的switch case以展示該方法的靈活性，突出CUDA如何有效利用模板參數來提升不同配置下的性能。

 // 添加這個函數來幫助展開并添加模板  
template <unsigned int blockSize>  
__device__ void warpReduce(volatile int* sdata, int tid){  if(blockSize >= 64) sdata[tid] += sdata[tid + 32];  if(blockSize >= 32) sdata[tid] += sdata[tid + 16];  if(blockSize >= 16) sdata[tid] += sdata[tid + 8];  if(blockSize >= 8) sdata[tid] += sdata[tid + 4];  if(blockSize >= 4) sdata[tid] += sdata[tid + 2];  if(blockSize >= 2) sdata[tid] += sdata[tid + 1];  
}  // 歸約 5 – 完全展開  
template <unsigned int blockSize>  
__global__ void reduce5(int *g_in_data, int *g_out_data){  extern __shared__ int sdata[];  // 存儲在共享內存中  // 每個線程從全局內存加載一個元素到共享內存  unsigned int tid = threadIdx.x;  unsigned int i = blockIdx.x*(blockDim.x*2) + threadIdx.x;  sdata[tid] = g_in_data[i] + g_in_data[i+blockDim.x];  __syncthreads();  // 分步執行歸約，減少線程同步  if (blockSize >= 512) {  if (tid < 256) { sdata[tid] += sdata[tid + 256]; } __syncthreads();  }  if (blockSize >= 256) {  if (tid < 128) { sdata[tid] += sdata[tid + 128]; } __syncthreads();  }  if (blockSize >= 128) {  if (tid < 64) { sdata[tid] += sdata[tid + 64]; } __syncthreads();  }  if (tid < 32) warpReduce<blockSize>(sdata, tid);  if (tid == 0){  g_out_data[blockIdx.x] = sdata[0];  }  }

相應地，需要修改內核調用方式以支持完全展開：

 // 完全展開所需  
// 啟動內核并同步線程  
switch (blockSize) {  case 512:  reduce6<512><<<num_blocks, 512, 512 * sizeof(int)>>>(dev_input_data, dev_output_data, n);  break;  case 256:  reduce6<256><<<num_blocks, 256, 256 * sizeof(int)>>>(dev_input_data, dev_output_data, n);  break;  case 128:  reduce6<128><<<num_blocks, 128, 128 * sizeof(int)>>>(dev_input_data, dev_output_data, n);  break;  }

實現方式與REDUCE-4相似，主要變化是將

blockSize

作為編譯時確定的模板參數。通過包含條件語句處理不同的

blockSize

值，以及使用switch語句根據這些值調用相應的內核版本。

結果

REDUCE-5 結果

靈活性與可擴展性的權衡

雖然Reduce5通過為已知塊大小完全展開循環來提升效率，但這種方法缺乏靈活性且難以擴展。完全展開技術嚴重依賴編譯時優化，將內核限制為固定的塊大小配置。當數據大小與塊配置不完全匹配時，可能導致效率降低和GPU資源的次優利用。此外，為每個塊大小管理多個內核版本增加了開發復雜度，限制了對變化工作負載的動態適應能力，使其在輸入大小變化較大的通用應用中實用性受限。

因此，需要借鑒REDUCE-3中加載時預歸約的思想，嘗試執行更多的加法操作而非僅限于第一次加法。

REDUCE-6：多重歸約與線程級并行

REDUCE-6通過引入"算法級聯"的動態方法來解決REDUCE-5中的剛性和可擴展性問題。該方法讓每個線程在更廣泛的塊大小范圍內執行多次加法操作，有效減少了對特定塊配置的依賴。這種靈活性使算法能夠更平滑地適應不同的數據規模，在更廣泛的場景中優化資源利用率。

通過結合順序和并行歸約技術，REDUCE-6最小化了延遲并最大化了吞吐量，特別適用于具有高內核啟動開銷和多樣化工作負載的環境。基于Brent定理的工作分配策略確保每個線程在整個歸約過程中都能以最優方式貢獻計算能力，在與硬件能力有效匹配的同時保持成本效率。

該方法的核心思想是每個線程在同步屏障之前處理多個元素對，從而在更多計算中攤銷同步成本，提升整體性能表現。

最終優化內核實現

 // 添加這個函數來幫助展開并添加模板  
template <unsigned int blockSize>  
__device__ void warpReduce(volatile int* sdata, unsigned int tid){  if(blockSize >= 64) sdata[tid] += sdata[tid + 32];  if(blockSize >= 32) sdata[tid] += sdata[tid + 16];  if(blockSize >= 16) sdata[tid] += sdata[tid + 8];  if(blockSize >= 8) sdata[tid] += sdata[tid + 4];  if(blockSize >= 4) sdata[tid] += sdata[tid + 2];  if(blockSize >= 2) sdata[tid] += sdata[tid + 1];  
}  // 歸約 6 – 多重加法/線程  
template <int blockSize>  
__global__ void reduce6(int *g_in_data, int *g_out_data, unsigned int n){  extern __shared__ int sdata[];  // 存儲在共享內存中  // 每個線程從全局內存加載一個元素到共享內存  unsigned int tid = threadIdx.x;  unsigned int i = blockIdx.x*(blockSize*2) + tid;  unsigned int gridSize = blockDim.x * 2 * gridDim.x;  sdata[tid] = 0;  while(i < n) {   sdata[tid] += g_in_data[i] + g_in_data[i + blockSize];   i += gridSize;   }  __syncthreads();  // 分步執行歸約，減少線程同步  if (blockSize >= 512) {  if (tid < 256) { sdata[tid] += sdata[tid + 256]; } __syncthreads();  }  if (blockSize >= 256) {  if (tid < 128) { sdata[tid] += sdata[tid + 128]; } __syncthreads();  }  if (blockSize >= 128) {  if (tid < 64) { sdata[tid] += sdata[tid + 64]; } __syncthreads();  }  if (tid < 32) warpReduce<blockSize>(sdata, tid);  if (tid == 0){  g_out_data[blockIdx.x] = sdata[0];  }  }

關鍵的技術創新體現在while循環中，每個線程直接在共享內存中執行多次加法操作。該循環設計為每次迭代聚合兩個數據元素，有效地將必要操作數量和與全局內存的交互頻次減半。線程從全局內存加載數據并添加到先前的累積值中，然后按總線程數的兩倍向前跳躍，確保在下一次迭代中處理另一對元素。這種模式顯著減少了每個線程在任何時刻需要處理的數據量，最大化了可用帶寬的利用率并最小化了訪問延遲。

性能評估與比較分析

REDUCE-6 結果

所有優化技術的性能對比

與NVIDIA基準的對比分析

本實現與NVIDIA官方實現的主要差異在于硬件平臺的不同。NVIDIA的研討會使用GeForce 8800，而本研究采用Tesla T4。由于Tesla T4具有更優化的架構，使得初始實現就具備了更好的性能基礎，但這也意味著性能提升的空間相對有限。雖然無法復現NVIDIA展示的戲劇性加速效果，但成功展示了持續的優化進程和GPU峰值性能的逐步提升。

總結

基于本次優化實踐，總結出CUDA內核優化的核心要點如下。

首先，深入理解核心性能特征是優化的基礎，包括內存合并訪問、分支發散管理、內存庫沖突解決以及延遲隱藏技術的應用。其次，充分利用性能指標進行瓶頸識別，通過計算和內存性能指標判斷內核是計算受限還是內存受限。第三，系統化地識別瓶頸來源，確定性能限制是由內存訪問、計算能力還是指令開銷造成的。第四，采用漸進式算法優化策略，先優化基礎算法邏輯，再進行循環展開等高級優化。最后，靈活運用模板參數技術進行代碼生成的精細調優，確保為不同塊大小配置提供最優的實現方案。

代碼：

https://avoid.overfit.cn/post/af59d0a6ce474b8fa7a8eafb2117a404

作者：Rimika Dhara