OpenCL Image 的妙用

曾提過許多的 OpenCL 的 wrapper 都不重視 Image 這個型別, 僅僅提供 Buffer 的物件, 這樣的作法除了錯誤地認為都僅只是 memory buffer 外, 忽略了 GPU 架構設計上的本質是相當可惜的.

在 CPU 上 memory access latency 的降低仰賴的是 cache, 對於 GPU 而言靠的是 wavefront(warp)/threads 的切換來吸收, 這裡以 Imagination 曾經介紹過其 PowerVR Rouge 的運作機制 來做說明 (其他的 GPU有作法不同但本質相同或類似的方式):

Rogue GPU 一個 Warp 基本上是 32-theads, 內部有著 16組 Warp slot 作排程切換, 由於每個 Warp 中的 32-threads 是以 lock-step 進行, 因此一旦有 thread 的執行狀態沒有被滿足整組 Warp 即被切為 idle state(紅色). 而如同上圖, GPU 是在期間一一切換尋找 ready state (黃色) 的 warp. 而如同影響 CPU 最大的一般, 其中最常出現的即為 memory latency. 一旦 memory latency 過高, 或是 bandwidth 不足就會發生下圖類似 CPU 的 memory stall 的情況:

那麼這與 Image的使用有何關連? - 因為多數的 GPU texture unit 有著自己的 memory pipeline 與 cache.  這表示以 texture unit 搭配既有的 memory pipeline 合併使用, 因此能夠相較於僅使用 Buffer 的情況下, 提供 PE/ALU 更高的 memory throughput.

像是 Qualcomm 提及(TP 即為為 Texture Pipe):

而像是 ARM Mali GPU 的說明也很清楚:

其他 Intel, NV, AMD GPU 都有著類似的方式.

在 cltk 中提供了一個簡單的範例 - gemm, 其中 gemm.cl 有兩個 kernel, 兩者主要差異在於一個完全使用 Buffer 而另一個有使用 Image. 相信很多人有疑問, 為何 GEMM 能夠以 Image 方式實作. 這仰賴兩件事

1. 以整數座標, 並且不使用 sampler 的方式來透過 read_imagef 來讀取 Image
2. data layout 的擺放方式為 CL_RGBA

滿足上述的方式, 就可以依序讀取 float4 (或是指定的資料型別)的向量.

使用 image 的 kernel:

__kernel void sgemm_buf_img(
   __global const float *A,
   const int lda,
   __global float *C,
   const int ldc,
   const int m,
   const int n,
   const int k,
   __read_only image2d_t Bi
){
    int gx = get_global_id(0) << 2;
    int gy = get_global_id(1) << 3;

    if ((gx < n) && (gy < m)){
        float4 a[8];
        float4 b[4];
        float4 c[8];

        for (int i = 0; i < 8; i++){
            c[i] = 0.0f;
        }
        A += gy * lda;

        for (int pos = 0; pos < k; pos += 4){
            for (int i = 0; i < 4; i++){
                b[i] = read_imagef(Bi, (int2)(gx >> 2, pos + i));
            }

            for (int i = 0; i < 8; i++){
                a[i] = vload4(0, A + mul24(i, lda) + pos);
                c[i] += a[i].x * b[0] + a[i].y * b[1] + a[i].z * b[2] + a[i].w * b[3];
            }

        }

        for (int i = 0; i < 8; i++){
            int C_offs = mul24((gy + i), ldc) + gx;
            vstore4(c[i], 0, C + C_offs);
        }
    }
}

不使用 Image 的 Kernel:

__kernel void sgemm_buf_only(
   __global const float *A,
   const int lda,
   __global float *C,
   const int ldc,
   const int m,
   const int n,
   const int k,
   __global const float *B
){

    int gx = get_global_id(0) << 2;
    int gy = get_global_id(1) << 3;

    if ((gx < n) && (gy < m)){
        float4 a[8];
        float4 b[4];
        float4 c[8];

        for (int i = 0; i < 8; i++){
            c[i] = 0.0f;
        }
        A += gy * lda;

        for (int pos = 0; pos < k; pos += 4){
            for (int i = 0; i < 4; i++){
                b[i] = vload4(0, B + (pos + i)*ldc + gx);
            }

            for (int i = 0; i < 8; i++){
                a[i] = vload4(0, A + mul24(i, lda) + pos);
                c[i] += a[i].x * b[0] + a[i].y * b[1] + a[i].z * b[2] + a[i].w * b[3];
            }

        }

        for (int i = 0; i < 8; i++){
            int C_offs = mul24((gy + i), ldc) + gx;
            vstore4(c[i], 0, C + C_offs);
        }
    }
}

以下為在個人的 Geforce 860M 上的執行結果

$ LD_LIBRARY_PATH=. ./gemm
build log:
cltk_image_alloc w[1024] h[1024] p[1024] us[16]
B pitch 1024 usize 16
pitch 16384
main test overhead gemm takes 103910 us
check result 1...
main test overhead gemm_buf takes 130500 us
check result 2...

兩者間有著近 30% 的效能差距, 若搭配 local memory 的使用或是在手機平台, 效能差距比例應會更明顯.

這個 GEMM 的範例主要參考自 Qualcomm 的 Optimization 文章, 並撰寫對比的 kernel

• Matrix Multiply on Adreno GPUs – Part 1: OpenCL Optimization
• Matrix Multiply on Adreno GPUs – Part 2: Host Code and Kernel