Qualcomm_Mobile_OpenCL.pdf 翻譯-6-工做組尺寸的性能優化

對於許多kernels來講，工做組大小的調整會是一種簡單有效的方法。這章將會介紹基於工做組大小的基礎知識，好比如何獲取工做組大小，爲何工做組大小很是重要，同時也會討論關於最優工做組大小的選擇和調整的通常方法。

 

6.1 獲取最大的工做組尺寸

         在運行完clBuildProgram後，使用下面的API函數能夠查詢設備的最大工做組尺寸。

       size_t maxWorkGroupSize;

       clGetKernelWorkGroupInfo(myKernel,

                                myDevice,

                                CL_KERNEL_WORK_GROUP_SIZE,

                                sizeof(size_t),

                                &maxWorkGroupSize,

                                NULL );

 

         在clEnqueueNDRangeKernel中使用的實際工做組尺寸不能超過maxWorkGroupSize。若是應用程序沒有指定工做組大小，Adreno OpenCL軟件可能會選擇最大的工做組尺寸。

        

6.2 須要的和優先的工做組尺寸

         每個kernel函數都有他須要或者優先的工做組大小。對於須要的工做組大小，OpenCL經過下面方法提供給編譯器。

• 使用reqd_work_group_size屬性。

        

         做爲需求，reqd_work_group_size(X, Y, Z) 屬性會傳入一個指定的工做組尺寸。若是指定的工做組大小不能知足將會返回一個錯誤。

         好比，若是要求16x16的工做組尺寸：

         __kernel __attribute__(( reqd_work_group_size(16, 16, 1) ))

       void myKernel( __global float4 *in, __global float4 *out)

       {  . . .  }

 

• 使用work_group_size_hint 屬性

 

         OpenCL會嘗試使用這個指定的尺寸，可是不保證真實的大小與指定的一致。好比，提示使用64x64工做組尺寸：

         __kernel __attribute__(( work_group_size_hint (64, 4, 1) ))

       void myKernel( __global float4 *in, __global float4 *out)

       {  . . .  }

 

         在許多狀況下，當工做組尺寸嚴格指定時，編譯器不能保證能編譯出最優的機器代碼。並且，若是片上寄存器不能知足要求的工做組尺寸時，編譯器可能會須要將寄存器溢出到系統的RAM內存上。所以，這兩種屬性並不建議使用，除非必須指定工做組尺寸，kernel才能運行。

         注意：爲了交叉編譯的兼容性，將kernel寫成依賴固定工做組的尺寸或者佈局，並非一個好的作好。

6.3 影響工做組最大尺寸的因素

         若是沒有指定工做組尺寸的屬性，一個kernel的最大工做組尺寸依賴如下的幾個因素：

• kernel的寄存器使用。通常來講，kernel越複雜，寄存器使用越多，支持的最大的工做組尺寸越小。過多地使用寄存器的緣由以下：
  • 　　一個工做項中有過多的工做任務。
  • 　　有控制流
  • 　　高精度的數學函數（好比，沒有使用內部函數或者快速數學運算的編譯選項-fastmath）
  • 　　本地內存，若是須要分配額外的寄存器暫時存儲裝載和存儲指令中源和目的地址。
  • 　　私有內存，好比爲每個工做組定義了一個數組
  • 　　循環展開
  • 　　內聯函數

        

• 通用寄存器的大小
  • 　　Adreno低級系列的GPU可能有更少的寄存器數量。

 

• n  kernel中的柵欄（Barrier）
  •         若是一個kernel沒有使用柵欄（barrier），在Adreno A4x and A5x系列中，在不用考慮寄存器使用的狀況下，工做組最大能夠設置爲DEVICEMAXIUMUM。

        

      6.4 沒有barrier的kernels

         之前地，一個workgroup中全部的work item要求在同一時間同時駐留在GPU上。對於大量消耗寄存器的kernel，這將會限制他們的最大工做組尺寸，並將會遠遠小於設備支持的最大工做組尺寸。

         從Adreno A4系列起，不須要考慮寄存器的使用狀況，沒有barrier的kernel就能夠有Adreno支持的最大工做組尺寸，通常是1024。對於這種類型的kernel（沒有barrier）來講，由於不須要wave之間進行同步，因此當一箇舊的wave執行完畢，新的wave就能夠開始執行了。

 

         在某些狀況下，擁有最大的workgroup尺寸並不意味着他們有最好的並行性。一個沒有barriers的kernel可能會由於太複雜致使只有不多的wave在SP上並行執行，這將會致使性能下降。開發者須要繼續優化和減小寄存器使用，不考慮從clGetKernelWorkGroupInfo函數中獲取到的最大的workgroup尺寸。

 

      6.5 工做組尺寸的調整

         這個部分將會介紹一些在選擇最優的工做組尺寸和形狀時通用的指導準則。

        

      6.5.1 避免使用默認的工做組大小

         若是一個kernel調用沒有指定workgroup的尺寸，那麼OpenCL會用簡單的方法找一個能用的工做組尺寸。開發者必需要要意識到，這種默認的工做組尺寸一般不是最優的。有效的作法是，手動嘗試使用不一樣的工做組大小和維度（2D/3D），而後找出最優的一個。

 

      6.5.2 越大的工做組尺寸，越好的性能？

         對於許多kernel來講這是正確的，由於增長工做組尺寸可以容許更多的wave運行在SP上，這樣可以更好地隱藏延遲和提高SP的使用。

         然而，對於某些kernel來講，增長工做組尺寸可能會致使性能退化。一種狀況是，因爲不良的數據局部性和訪問模式，越大的工做組尺寸將致使越多的cache垃圾。這個數據局部性的問題在使用texture獲取時更加嚴重，由於texture cache比統一的L2 cache要小。最終，決定最優的工做大小和維度的本質是kernel的數據獲取。

 

6.5.3 固定的 vs. 動態的工做組尺寸

         爲了避免同設備之間的性能兼容性，避免假設一個工做組尺寸可以適合全部的設備，避免對workgroup尺寸固定編碼。一個指定的工做組大小和維度在一個設備上是最優的，在另外一個設備上多是次優的。所以，給定一個kernel，建議針對kernel可以執行的全部設備統計出不一樣的workgroup尺寸，而後在運行時對每一個設備選出一個最優的。

 

6.5.1 一維 vs二維 vs三維 workgroup（1D/2D/3D）

         kernel的維度能夠會影響性能。取決於work item的數據獲取方式，在某些狀況下，一個2D的kernel可能會在cache上有更好的數據本地性（數據在cache上），致使更好的內存獲取和更好的性能。然而在其餘狀況下，一個2Dkernel比1D會產生更多的cache 垃圾。建議嘗試使用不一樣的維度，從而獲取最優的性能。

 

6.6 關於workgroup的其餘話題

      6.6.1 全局的work size和填充

         OpenCL 1.x 要求一個kernel的全局worksize 必須是workgroup尺寸的倍數。若是應用程序指定的workgroup 尺寸不知足這個條件，那麼clEnqueueNDRangeKernel的函數調用將會返回一個錯誤。在這種狀況下，應用程序能夠填充全局worksize，保證它是用戶指定的workgroup尺寸的倍數。

 

         注意：OpenCL 2.0 取消的這個限制，並且global worksize 並不須要必須是workgroup size的倍數，這種被叫作非歸一化的workgroup。

 

         理想狀況是，workgroup 尺寸的第一個維度是wave尺寸的倍數（好比說32），這樣能充分利用wave的資源。若是不是這種狀況，能夠考慮填充workgroup的大小來知足這個條件，須要記住，在OpenCL 1.x中，全局的worksize必須填充（保證是workgroup的倍數）。

      6.6.2 殘酷地尋找

         由於workgroup 尺寸選擇的複雜性，經驗經常是發現最優大小和維度的最好方法。

        

         一種選擇是，在程序開始時，使用一個與實際的工做任務相同複雜度（可是通常使用比較簡單的任務）的喚醒功能的kernel 去動態的尋找最優的workgroup 尺寸。而後將這個選出來的workgroup尺寸用在實際的kernel中。不少商業的標準檢查程序就是使用的這種方法。

      6.6.3 在workgroup中避免不均勻的工做負載

         一些應用程序可能被寫成，在不一樣工做組中出現不均衡的負載。好比說，基於區域的圖像處理的用例中會出現一些區域須要比其餘區域多不少處理的狀況。這種狀況須要避免，由於這會致使性能的不可預測性。另外的，若是單個workgroup任務須要太長時間運行的話，會致使上下文切換變的複雜。

         解決這個問題的方法是，使用兩個階段處理策略。第一個階段可能會收集感興趣的點和爲第二階段準備數據。工做負載越具備肯定性，在不一樣workgroups中進行均等的分配將會更簡單。

      6.6.4 工做組的同步

         OpenCL 並不能保證workgroup的執行順序，並且也沒有定義一個工做組同步的機制。不建議有須要依賴工做組順序的程序。

 

         在實際狀況下，可使用atomic 函數或者其餘方法，在workgroup之間能夠進行有限的同步。好比說，一個應用程序可能分配了一個全局內存對象，這個對象須要被不一樣的工做組中的workitem更新。一個workgroup能夠管理一個由其餘workgroup更新的內存對象。經過這種方式，可能會實現有限的工做組同步。