2021年9月6日 星期一

GPU 架構 淺談 : IMR, TBR, TBDR

        最近工作回歸到 GPU 領域, 也再次興起寫些東西的想法, 回顧blog 的歷史紀錄,發現上次寫文章, 已經是 8 年前的事了, 心中除了些許震驚之外, 也隱隱感嘆歲月流逝是如此的迅速. 

在這段時間, 整個 GPU 或是 Rendering 的技術推陳出新. 有需多的技術出現了. 有些是新的概念, 有些則是舊有的想法逐漸演進, 隨著硬體環境以及標準的演進與成熟, 而得以實現, 因此可以說隨手拾起都是新東西,例如 Vulkan , NV(Turing), Imagination (GR6500) GPU 的 Ray Tracing.

所謂 千里之路 始於足下, 縱使有這麼多的新技術與想法可以讓我們可以去學習, 我們還是要從邁開腳下的第一步開始. 因此我想就用 GPU 基本架構分類作為我再次出發時邁出的第一步.


讓我們開始吧.

目前三種 GPU 架構, 分別是 IMR, TBR 以及 TBDR. 

IMR : Immediate Mode Rendering

這應該是最直覺且容易理解的一種 rendering path, IMR 的 block diagram 可以參考圖一


         IMR 架構的 GPU 處理 primitive 的先後順序是根據 Primitive 在 Vertex Data 中的順序以及調用 Drawcall 的先後順序決定的.例如 Drawcall #1 送出 Triangle A  然後是 Triangle B, 接著 Drawcall #2 送出 Triangle C, 則 IMR GPU 處理 Triangle 的順序依序是 A, B 然後是 C. 

         因為這樣的特性, 一般來說 IMR 架構相比其他兩種架構會有較高的帶寬(memory bandwidth)用量 以及 overdraw(同一個位置的pixel 被重複寫入多次) , 為什麼會這樣呢?
在說明原因之前, 我們先看一下 IMR 架構的 GPU 處理 primitive 的 流程:
         
以前面的 Drawcall #1 為例子, GPU 會先處理 Triangle A, Vertex Processing 階段會從 memory 提取 Vertex Data, 進行座標轉換, 在下一級進行 view plane Clipping 以及 backface  Culling.

Rasterization 會將 geometry data  轉換成 fragment 並內插出per-fragment 的 attributes(e.g. depth, color, texture coordinate,..), 如果 fragment shader 裡面不會調整新的 depth 值則 early depth testing 會先進行 depth testing 以減少進入 fragment processing 的工作量. 否則就 bypass early depth testing.

進入 fragment processing, 每個 fragment 會從 memory 提取 texture data 並進行 fragment 計算(shading), 如果 fragment shader 會調整 depth 則接著下一步到 later depth testing 進行 depth testing, 這裡可能會需要再讀寫 Depth buffer. 

最後到 ROP, 如果需要 alpha blending, 還會需要讀取memory(framebuffer), 進行 alpha blending 計算後再寫回 memory.



圖二是處理完  Triangle A後, framebuffer 的內容.
歸納以上的描述可知, Triangle A 的每個 fragment 都會占用帶寬, 以及GPU 的計算量. 

接著, GPU 使用相同的流程處理 Drawcall #1 的Triangle B. 我們假設 Triangle B 的每個 fragment 的 depth 都是小於 Triangle A 的, 也就是說, 整個 Triangle B 距離 view point 的距離小於 Triangle A.

GPU 處理完 Triangle B 後, Framebuffer 的內容如圖三.

從圖三可以看出, 原本 triangle A 所涵蓋的 pixel, 有一部分被後來的 Triangle B 覆蓋了
這些被覆蓋的 pixel 最後不會出現在 framebuffer 中, 因此在 Triangle A階段, 這些 fragment 所使用的帶寬以及計算量都被浪費了,  這種同一個 pixel 被多次複寫的狀況稱做 overdraw. 

在圖四中用橘色標出 overdraw 的部分. 在這個例子裡 Triangle A 大約 84% 的帶寬用量和計算量因為被 Triangle B  overdrawing 而浪費了.


PC 平台的 GPU 以 IMR 架構為主, 在 PC 平台,  GPU 會有自己專用的 local memory, 不會
共用系統的 system memory,  因為不會搶到 系統上其它模組的帶寬, 因此即使多消耗帶寬, 對系統的影響不大. 且 IMR GPU 架構相對簡單, 除了上述的開銷, 沒有其它架構所特有的 overhead.

雖然 IMR 的額外開銷在 PC 平台上可以被接受, 但是 mobile 平台狀況並不相同, 在 mobile 平台上, GPU 和其它 SOC 上的模組共用 system memory, 只是把 system memory 的一個區塊劃分給 GPU 使用. 帶寬的使用量和使用者體驗會有關聯性.

如果 GPU 使用的帶寬過高, 會影響到 SOC 中的其它 module 可用的帶寬量(每個 module 都有優先權高低之分), 而且 帶寬的使用量越多, 代表會有越高的功耗, 在 mobile 平台上, 如何減少功耗是一個重要的課題, 越低的功耗, 除了代表越長的使用時間, 也會減少發熱, 產生的熱越少, 則發生 thermal throttle 導致系統的效能降低的機會越少, 就會有比較良好穩定的使用者體驗.


目前在 mobile 平台上的 GPU 架構有兩種, 分別是 TBR, 以及 TBDR.

TBR : Tile Based Rendering
圖五是 TBR GPU 的架構方塊圖.


Tile Based Rendering, 顧名思義, 是將整個 screen 分割成大小相同的區塊(e.g. 16 x 16 pixels), 區塊稱為 tile.

在一個 frame 中應用程式調用 drawcall 時, 會先由 tiling 處理每個 drawcall 的 primitives, 標定有哪些 title 被 primitives 覆蓋到, 把 primitive 的 id 記錄到該 tile 所屬 primitive list 中. 下圖以前面的 drawcall #1 為例說明 tiling.



app 調用完所有 drawcall 後, 最後會調用 eglswapbuffer (以 OpenGL/OpenGLES 為例) 完成一張 frame, 此時 GPU 會開始進行 rendering 階段, GPU 依次處裡每個 tile. 將 tile 中所有 primitives 處理完後, 再處理下一個 tile 的所有 primitives. 

因為tile size (e.g. 16 x 16) 遠小於 screen size (e.g. 720P), 因此 tile 上 shading 運算時所需的 memory 可以使用 on-chip memory, 因此不管 per-tile shading 的overdraw 有多高, 所有的 memory r/w 都不需要出到SOC 外部的 system memory, 只需在 per-tile 的所有primitives 都完成 shading 後, 一次性將 tile 的color buffer 寫到 frame buffer. 
因為使用了 on-chip memory, 因此減少功耗以及system memory 帶寬使用量. 這讓 TBR 架構比 IMR 架構更適合適用在 mobile 平台. 目前 ARM 的 mali GPU 就是屬於 TBR 架構 GPU.

從前面的描述可知, TBR 減少了系統帶寬的使用量, 但是每個 tile 中的 primitive 繪製還是會引發 overdraw, 這會讓使用了 GPU 計算資源的 fragment , 可能被後來的其它 primitive 的 fragment 複寫而導致計算浪費, 所以overdraw 的問題沒有得到解決. 因此有了 TBDR (Tile Based Deferred Rendering) 架構.

TBDR

TBDR 除了有 tiling 這個 "deferred", 另外 HSR & tag buffer, 是另一個 "deferred" 的部分. HSR (Hidden Surface Remove) 的目的是消除 tile 中的 overdraw, HSR 會遍歷 tile 中所有的 primitives.找出最後真正會被畫出的 fragment. 並用 tag buffer 記錄這些 fragment 所屬的 primitive.
只有這些真正visible 的 fragment 會進行 fragment processing. TBDR 架構可以真正做到 零 overdraw 並與場景 drawcall 順序完全無關, 這可以最大限度的節省 fragment processing 的運算資源. 獲得的好處基本上會超過 "deferred" 處裡本身須付出的 overhead. 目前 Imagination 的 GPU 就是採用 TBDR 架構.


2013年12月11日 星期三

OpenCL 模擬 iOS7 毛玻璃效果(frosted glass)

最近筆者將手機作業系統升級到 iOS7, 發現 iOS7 的介面有一個亮點 : 毛玻璃(frosted glass)效果, 透過毛玻璃所在的 layer, 可以隱約看到下層 layer 的內容. 如圖一所示.

圖一
搜尋一些資訊後發現 不少人在討論這個效果, 因為 apple 使用的方法無法確切得知, 所以筆者就以自己的想法實作.  首先 需要一個用以柔化影像的 low-pass filter. 筆者計畫使用 Gaussian filter 並且使用 GPU 做 Gaussian filter convolution. 關於 Gaussian filter, 可以參考下面兩個在 wiki 的介紹:
Gaussian kernel 的大小由 sigma 決定, Host 端的程式根據 sigma 決定 kernel size, 根據 gaussian filter 的定義計算 mask 中的每個常數, 將 mask 儲存在 OpenCL 的 image object.    kernel size 以及 mask 會給 OpenCL kernel 進行 convolution 時使用. 用以計算 Gaussian mask 的程式片段如下:

void
CreateBlurMask(float sigma, int * maskSizePointer) 
{
int maskSize = (int)ceil(3.0f*sigma);
g_mask.resize((maskSize*2+1)*(maskSize*2+1));

float sum = 0.0f;
float* pmask = &g_mask[0];
pmask += maskSize*(maskSize*2+1) + maskSize;

for(int a = -maskSize; a < maskSize+1; a++)
{
for(int b = -maskSize; b < maskSize+1; b++) 
{
float temp = exp(-((float)(a*a+b*b) / (2*sigma*sigma)));
sum += temp;
pmask[a+b*(maskSize*2+1)] = temp;
}
}
// Normalize the mask
for(int i = 0; i < (maskSize*2+1)*(maskSize*2+1); i++)
g_mask[i] = g_mask[i] / sum;

*maskSizePointer = maskSize;
}

根據我前一篇文章 OpenCL介紹 作修改, 在 host 端程式代碼的初始化, 大致上增加了以下的部分.
調用 CreateBlurMask, 並且給定 sigma = 3.0

CreateBlurMask(3.0, &g_maskSize);

接著, 建立一個 image object, 並且將 Gaussian mask 的內容由 host 端的 memory 複製到 device 端的 memory:
imgObject[2] = clCreateImage2D(GPU_context, CL_MEM_READ_ONLY,                                                                  &mask_fmt, g_maskSize*2+1, 
                                                       g_maskSize*2+1, 0, NULL, NULL);

err = clEnqueueWriteImage(commandQueue, imgObject[2], CL_TRUE, origin, 
                                         mask_region, 0, 0, &g_mask[0], 0, NULL, NULL);

if(err != CL_SUCCESS) 
{
cleanUpOpenCL(GPU_context, commandQueue, program, kernel, 
         memoryObjects, numberOfMemoryObjects,
 imgObject, numberOfImgObjects);
return 0;

}

設定好送進 kernel function 的參數
clSetKernelArg(kernel, 2, sizeof(cl_mem), &imgObject[2]);  // source image
clSetKernelArg(kernel, 3, sizeof(cl_int), &g_maskSize);       //  mask Size


Host 端處理每個 frame 的流程如下:
(1) 產生一份原始影像 scaling down 成較低解析度後的結果.
(2) 讓GPU對 較低解析度的影像 進行 convolution.
(3) 讀回結果.
(4) CPU 將讀回的結果scaling up 成和原始影像一樣解析度.
(5) CPU 根據 ROI,  將scaling up後的影像和原始影像合成.

圖二

因為要 composition, 故需要產生一份 composition 使用的 mask. 此mask 在 ROI 內的值為 0xff,其餘部分皆為 0.

camera >> Img1;
Mat mask(Img1.rows, Img1.cols, CV_8UC1);
mask.setTo(0);

Mat roi(mask, Rect(140,100,340,260));
roi = 0xff;

composition 的代碼如下, ocl_img 是從讀回讀回並 scaling up 的結果. Img0 則是原始影像.
ocl_img *= 0.3;
ocl_img.copyTo( Img0, mask );

接下來看 Kernel Function 的內容:
__constant sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE |
                               CLK_ADDRESS_CLAMP | 
                               CLK_FILTER_NEAREST;

__kernel void laplacian( read_only  image2d_t src_image,                         
                         write_only image2d_t dst_image,
                         read_only  image2d_t mask,
                         __private int maskSize )
{
int x = get_global_id(0);
int y = get_global_id(1);

int2 pos = {get_global_id(0), get_global_id(1)};


    // Collect neighbor values and multiply with Gaussian
    float4 fSum  = 0.0f;
    float4 fMask;
    
    float4 fColor;
    uint4  uiColor;    
    
    for(int a = -maskSize; a < maskSize+1; a++) 
    {
        for(int b = -maskSize; b < maskSize+1; b++) 
        {
            int2 coord = (int2)((a+maskSize, b+maskSize));
            uiColor = read_imageui(src_image, sampler, pos + (int2)(a,b));
            fColor  = convert_float4(uiColor);
            
            fMask   = read_imagef(mask, sampler, coord);
            
            fSum += fColor * fMask.x;
        }
    }
       
    uiColor = convert_uint4 (fSum);   
    write_imageui(dst_image, pos, uiColor);    
}
Kernel 這邊的代碼 根據 mask size 進行 convolution, 再將結果寫到 output image object.

2013年11月12日 星期二

OpenCL 介紹

前言

在以往只有 fixed function pipeline 的時代, GPU 單純做為 3D 繪圖加速器, 應用程式透過諸如 OpenGL, OpenGL ES, Direct3D 等等的 3D API, 使用硬體提供的繪圖加速服務, 在 programmable pipeline 出現之後, GPU的應用進入了另一個領域 : computing. 這個階段出現了一個新的名詞 : GPGPU; General Purpose Computing on Graphics Processing Unit. GPU 內的 shader, texturing pipeline, arithmetic pipeline以及 Multi-thread 架構, 開始被應用在計算領域. 在這個階段, 計算應用必須使用標準的 3D API來完成, 這對計算應用在實作上造成不便, 例如, 要實做一個簡單的浮點數資料相加計算程式, 需要涉及到像是 vertex data, texture 等 和計算無關的資料, 概念和流程.
為了讓 GPU 更方便使用在計算領域, 專門為 computing 而制定的 API 出現了,
這讓 GPU 在 computing 的應用進入了第二個階段, 這個階段的概念是 Heterogeneous Computing; 異質性計算, 要能使用系統上不同性質的運算單元, 舉例來說, 在 CPU與 GPU 的協同運算中, 由CPU 負責程式的流程控制, 大量的資料則交由 GPU 以平行化的方式進行計算.
常見的應用包括:
image processing
computer vision
speech processing

可以參考筆者另一篇文章 OpenCL 毛玻璃效果

目前業界主要的 compute API 有以下四種:
(1) OpenCL ; Apple 首先提出, 並聯合多家廠商合組 working group, 目前由 Khronos 維護. 是本文章要介紹的API.
(2)C++AMP, DirectCompute ; 由 MicroSoft 主導制定.
(3)CUDA ; nVidia 自家的規格.
(4)Renderscript ; Google 制定, Android 3.0 開始出現.

其實使用這些 API 的目的不外乎就是要藉 GPU 的平行運算能力, 縮短計算任務所需的時間.
因此以平行化運算的方式思考, 應該是使用 GPU 加速運算的重點.
Laplace filter 影像處理














2013年8月23日 星期五

在程式中開關 log 的方法

撰寫程式, 不論是為了除錯, 或是了解程式運作的狀態, 需要在程式代碼中 加進一些 log.
一旦加入 log, 當然也需要一個開啟和關閉 log 的機制. 而且必須是一個方便使用的開關機制.

一般來說, 我們可能會使用 以下的方式在程式代碼中加入 log

  #define MY_DEBUG
  .......

  #ifdef  MY_DEBUG 
  printf("This is a log.\n");
  #endif
  .......

使用 MY_DEBUG 控制 log 機制的開啟和關閉.
這種做法的缺點是 加入 log 的同時也必須跟著加入 #ifdef 和 #endif. 撰寫上不方便 且程式代碼會看起來比較凌亂.

2013年8月15日 星期四

Linux PID 何時會被重複使用

在 Android 中, 大部分的 app在執行時, 若是使用者按下 back 鍵 回到桌面, 該 app 的 process 會暫停執行, 但是 process 可能仍然會存在系統中. 如圖一所示是尚未執行 app 前的系統 process 狀態:
圖一 未執行app 之前系統中的 process

2013年2月5日 星期二

OpenGL projection matrix

OpenGL projection matrix

      

       為了將 3D 場景呈現在 2D 的顯示裝置, 必須透過 projection transform 3D 座標 轉換成 2D 座標, 不論是 OpenGL ES 1.1 fixed function pipeline或是 OpenGL ES 2.0/3.0 vertex shader, 都需要使用投影矩陣(projection matrix)來進行 projection transform. 根據投影方式Projection matrix 分為兩種, 分別是透視投影 (perspective projection) 以及平行投影 (parallel projection), 本文說明如何推導 OpenGL ES perspective projection matrix.