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.


2013年1月29日 星期二

攝像頭校正 camera calibration - part3 OpenCV programming

攝像頭校正 camera calibration - part3 OpenCV programming

在前兩篇文章  camera calibration -part1 camera model 以及 camera calibration -part2 camera calibration 中介紹了 camera 內部參數(intrinsic parameter) , 形變參數 (distortion parameter) 以及校正的方法和原理, 本文說明如何使用 OpenCV 實作camera校正 程序.
校正的程序分為兩部分;


2012年12月16日 星期日

Histogram Equalization 直方圖等化

Histogram Equalization

當透過 sensor 抓取影像, 我們可能會需要提高影像的品質, 使影像的細節或特徵更為凸顯. 其中一個方法是對影像亮度的分佈進行改變, 此類方法稱為 對比增強. 本文所要介紹的 Cumulative Histogram Equalization(直方圖等化) 法即是影像增強的技術之一.

直方圖等化使用機率分佈, 將原先的亮度分佈重新均勻的等化到新的亮度值. 原來在直方圖分佈中 比較窄化的影像, 例如 : 低對比度的影像, 透過直方圖等化的過程, 其亮度灰階的分佈變成均勻分散, 而成為高對比度的影像. 另外, 對於影像偏暗的部分提高亮度, 偏亮的部分則降低亮度, 使得細節呈現更為清晰.


Cumulative Histogram Equalization 的步驟如下
1.     建立影像的 histogram(PDF).
2.     計算影像的 cumulative distribution function(CDF).
3.     根據 CDF 以及 cumulative equalization 公式 計算灰階亮度的對應關係.
4.     根據對應關係計算出新的灰階亮度.


2012年11月15日 星期四

skia 性能優化 - part 2


Android 在 3.0 版之前, 包括Android 2.2, 原生UI 的效能一直讓使用者有所微詞. 部分的原因在於UI 上的所有圖形, 都使用 CPU rendering. 要提高 UI  rendering 的效能, 除了改以特定硬體加速, 例如 GPU 或是 2D engine, 另外一個方法就是針對程式代碼進行優化, 包括針對效能
瓶頸部分, 改用處理器所提供更有效率的指令完成. 當然, 繪圖的加速只是一部分, 另外像是 


2012年10月23日 星期二

攝像頭校正 camera calibration - part 2 calibration


Lens Distortion

在前一篇文章 攝像頭校正 camera calibration - part 1 camera model提到 使用 lens 的好處是可以增加進光量, 縮短曝光時間. 但是缺點則是會造成非線性的影像形變.

常見的兩種主要形變, 分別是徑向形變( radial distortion)以及 切向形變 (tangent distortion). 徑向形變是由於 lens 的形狀, 使得距離lens中心越遠的光線在通過 lens時彎折較大, 越靠近 lens 中心的光線則彎折越小. 切向形變則是由 camera 模組組裝的誤差造成的. 主要是 lens 和 sensor 並非平行對齊. 圖一是兩種形變的影像.


2012年10月9日 星期二

skia 性能優化


Skia是一個 Android 中的一個 open source 向量圖型庫,

在 Android 2.2 中 skia 幾乎無所不在, 提供了 draw text, draw rectangle, draw bitmap 等等的功能. 許多 apk 的UI, 其具體實現就是用skia, 例如, launcher. 而 skia 在 Android 的原生代碼中, 提供了 skia 以及 skiaGL (OpenGL加速的版本). 在 skia 的版本中,  skia 的運算仰賴 CPU, 雖然有提供了 fixed point 以及 floating point 兩套版本, 但筆者實際上透過 #define 調整時, 編譯上會產生錯誤. 有部分的代碼需要修改調整.(是否有其他同好有調試過這部分的代碼?)


2012年9月14日 星期五

擴增實境 - A simple Augmented Reality demo


A simple Augmented Reality demo

最近研究了一個有趣的主題 : Augmented Reality ; AR, 中文翻譯為 擴增實境. 並且用目前的研究心得寫了一個簡單的 demo program.

這是一個 使用 marker 的 AR demo program. 其原理是使用 webcam 抓取影像, 對影像做 marker detection, 根據四邊形的 marker 在 2D image 上的形變狀況 推導出 marker 所在之 3D 平面的 3D transform. 然後使用此 3D transform, 在抓取的影像中繪制 3D 物件. 如此, 就把真實的影像 和虛擬的物件融合在一起.


2012年9月8日 星期六

攝像頭校正 camera calibration - part 1 camera model

PART 1 Camera Model


所謂 Camera model 是使用幾何(geometry)的方式描述真實世界中的物體在 camera 影像平面上的成像方式. 本文會說明最基本的 camera model, 也就是針孔攝像頭模型(pinhole camera model). 另外, 使用鏡片(lens)可以增加進入成像平面的單位時間光線總量,因此縮短了曝光所需要的時間, 但是使用lens也造成了影像或多或少產生了畸變(distortions), 因此我們還需要了解 lens distortion model.
有了camera 的 geometry model 以及 lens 的 distortion model 就可以建構完整的 camera 的內部參數(intrinsic parameter),   camera calibration 需要使用到 camera 的內部參數.


2012年9月1日 星期六

Android 建立 surface 的過程


Android 的圖形系統是 Client-Server 架構, Server 是 SurfaceFlinger, Client 則是應用程式.
應用程式第一次執行時, Window Management Service (WMS) 會開啟一個 session, openSession()函式會被調用. openSession()函式在BSP的位置和程式代碼如下:


2012年8月14日 星期二

探討 Android gralloc HAL – part2

gralloc HAL有幾個比較重要的函式, 分別是 gralloc_alloc, gralloc_alloc_buffer,  init_pmem_area,  gralloc_free, 另外, 還有一個class : SimpleBestFitAllocator.
這些函式的程式代碼在

2012年8月2日 星期四

探討 Android gralloc HAL – part 1


整個 Android 2.2的圖型模組, 不論是 2D, 3D應用程式運行, 只要是圖形相關的應用, 都需要配置 graphic buffer, 在 Android 中負責這工作的模組就是 gralloc HAL(gralloc 硬體抽象層).

Gralloc HAL 的程式碼在 /hardware/libhardware/modules/gralloc/Gralloc.cpp
動態連結庫 gralloc.default.so 位置在系統路徑 /system/lib/hw/


2012年8月1日 星期三

Android – screen compatibility mode


工程師的工作不外乎要 debug, 最近又遇到一個 issue.

在 WVGA(width=800, height=480) 的平台上, 執行應用程式, 該應用程式卻 allocate(width=533, height=320)的 graphic buffer. 造成畫面內容顯示不正確.


2012年7月29日 星期日

Android 系統語系設定


最近因為要解一個 Android 平台的 issue 而接觸到了 Android 系統的語系, Android 的語系設定主要有兩個部分

系統初設(default) 的語系設定
透過 settings 改變語系設定.

系統 初設的語系設定

2012年7月21日 星期六

OpenCV 影像擷取 - 使用 DirectShow


在本篇文章中, 筆者將說明如何使用 DirectShow  擷取影像. 當初使用 DirectShow 是因為
需要從事 Stereo vision相關的研究, 要從多個 camera 擷取影像, 且使用 OpenCV 2.2版, 但是這一版的 OpenCV 原生函式無法正確的擷取多個 camera 的影像. 因此才改用 DirectShow.

2012年7月14日 星期六

OpenCV 影像擷取


本文介紹如何使用 OpenCV 進行影像擷取,  在前一篇文章中 Hello_OpenCV 從檔案將圖檔讀入並且顯示, 現在, 我們進一步擴充 Hello_OpenCV, 把原先的從檔案讀取圖檔改成從 WebCAM 擷取影像.
OpenCV 提供的影像擷取函式有兩個, 分別是
                               CvCapture* cvCaptureFromCAM(int index);
以及
                               IplImage* cvQueryFrame( CvCapture* capture );