cuda bilinear 紋理記憶體、共享記憶體實作與解說

之前介紹過一篇

• [原始碼] C / C++ 線性插補 bilinear 與注意事項

裡面詳細介紹注意事項與如何實作，大致簡單說一下就是等比取值

如果 A=0, B=10 ，那麼 AB=1 大概就是這個原理，如果是二維平面4個點就是取3次剛剛的過程，先兩次兩個點，然後再把兩個點的結果取一次。

CPU函式

上面那個敘述就是我們的主要函式了，大致可以寫成如下程式

__host__ __device__
inline static float bilinearRead(const float* img, 
    size_t width, float y, float x) // 線性取值
{
    // 獲取鄰點(不能用 1+)
    size_t x0 = floor(x);
    size_t x1 = ceil(x);
    size_t y0 = floor(y);
    size_t y1 = ceil(y);
    // 獲取比例(只能用 1-)
    float dx1 = x - x0;
    float dx2 = 1 - dx1;
    float dy1 = y - y0;
    float dy2 = 1 - dy1;
    // 獲取點
    const float& A = img[y0*width + x0];
    const float& B = img[y0*width + x1];
    const float& C = img[y1*width + x0];
    const float& D = img[y1*width + x1];
    // 乘出比例(要交叉)
    float AB = A*dx2 + B*dx1;
    float CD = C*dx2 + D*dx1;
    float X = AB*dy2 + CD*dy1;
    return X;
}

要操作這個函式大致上就用兩個for用x, y跑過圖中所有的點就可以了

__host__ void biliner_CPU_core(vector<float>& img, const vector<float>& img_ori, 
    size_t width, size_t height, float Ratio)
{
    int newH = static_cast<int>(floor(height * Ratio));
    int newW = static_cast<int>(floor(width  * Ratio));
    img.resize(newH*newW);
    // 跑新圖座標
    for (int j = 0; j < newH; ++j) {
        for (int i = 0; i < newW; ++i) {
            // 調整對齊
            float srcY, srcX;
            if (Ratio < 1) {
                srcY = ((j+0.5f)/Ratio) - 0.5;
                srcX = ((i+0.5f)/Ratio) - 0.5;
            } else {
                srcY = j * (height-1.f) / (newH-1.f);
                srcX = i * (width -1.f) / (newW-1.f);
            }
            // 獲取插補值
            img[j*newW + i] = bilinearRead(img_ori.data(), width, srcY, srcX);
        }
    }
}

共享記憶體

共享記憶體這邊有引用到之前寫的函式庫，詳細可以參考這篇站內文
https://charlottehong.blogspot.tw/2018/01/cuda.html

// 共享記憶體線性取值核心
__global__ void biliner_share_kernel(float* dst, const float* src, int srcW, int srcH, float ratio) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    int j = blockIdx.y * blockDim.y + threadIdx.y;
    int newH = (int)(floor(srcH * ratio));
    int newW = (int)(floor(srcW * ratio));
    if(i < srcW*ratio && j < srcH*ratio) { // 會多跑一點點要擋掉
        // 調整對齊
        float srcY, srcX;
        if (ratio < 1) {
            srcY = ((j+0.5f)/ratio) - 0.5;
            srcX = ((i+0.5f)/ratio) - 0.5;
        } else {
            srcY = j * (srcH-1.f) / (newH-1.f);
            srcX = i * (srcW -1.f) / (newW-1.f);
        }
        // 獲取插補值
        dst[j*newW + i] = bilinearRead(src, srcW, srcY, srcX);
    }
}

差不多就是從CPU那個函式改過來的而已，使用時候要先new兩個空間

// 共享記憶體線性取值函式
__host__ void biliner_share_core(float *dst, const float* src,
    size_t srcW, size_t srcH, float ratio)
{
    Timer T; T.priSta = 1;
    // 設置GPU所需長度
    int srcSize = srcW*srcH;
    int dstSize = srcSize*ratio*ratio;

    // 要求GPU空間
    T.start();
    CudaData<float> gpu_src(srcSize);
    T.print("  GPU new 空間1");
    T.start();
    CudaData<float> gpu_dst(dstSize);
    T.print("  GPU new 空間2");
    // 複製到GPU
    T.start();
    gpu_src.memcpyIn(src, srcSize);
    T.print("  GPU 複製");

    // 設置執行緒
    dim3 block(BLOCK_DIM, BLOCK_DIM);
    dim3 grid(ceil((float)srcW*ratio / BLOCK_DIM), ceil((float)srcH*ratio / BLOCK_DIM));
    T.start();
    biliner_share_kernel <<< grid, block >> > (gpu_dst, gpu_src, srcW, srcH, ratio);
    T.print("  核心計算");

    // 取出GPU值
    T.start();
    gpu_dst.memcpyOut(dst, dstSize);
    T.print("  GPU 取出資料");

    // 釋放GPU空間
    T.start();
    gpu_src.~CudaData();
    gpu_dst.~CudaData();
    T.print("  GPU 釋放空間");
}

紋理記憶體

這個速度最快代碼也最少，可以直接調用cuda內優化過的方法

// 宣告GPU紋理變數(只能放全域)
texture<float, 2, cudaReadModeElementType> rT;
// 紋理線性取值核心
__global__ void biliner_texture_kernel(float* dst, int srcW, int srcH, float ratio) {
    int idxX = blockIdx.x * blockDim.x + threadIdx.x,
        idxY = blockIdx.y * blockDim.y + threadIdx.y;
    if(idxX < srcW*ratio && idxY < srcH*ratio) { // 會多跑一點點要擋掉
        float srcX = idxX / ratio;
        float srcY = idxY / ratio;
        size_t idx = (idxY*srcW*ratio + idxX);
        dst[idx] = tex2D(rT, srcX+0.5, srcY+0.5);
    }
}

使用

__host__ void biliner_texture_core(float *dst, const float* src,
    size_t dstW, size_t dstH, float ratio)
{
    Timer T; T.priSta = 1;
    // 設置GPU所需長度
    int srcSize = dstW*dstH;
    int dstSize = srcSize*ratio*ratio;

    // 宣告 texture2D陣列並綁定
    T.start();
    CudaMemArr<float> cuArray(src, dstW, dstH);
    cudaBindTextureToArray(rT, cuArray);
    T.print("  GPU new 紋理空間+複製");

    // 設置 插植模式and超出邊界補邊界
    rT.filterMode = cudaFilterModeLinear;
    rT.addressMode[0] = cudaAddressModeClamp;
    rT.addressMode[1] = cudaAddressModeClamp;

    // 要求GPU空間
    T.start();
    CudaData<float> gpu_dst(dstSize);
    T.print("  GPU new 一般空間");

    // 設置執行緒
    dim3 block(BLOCK_DIM, BLOCK_DIM);
    dim3 grid(ceil((float)dstW*ratio / BLOCK_DIM), ceil((float)dstH*ratio / BLOCK_DIM));
    T.start();
    biliner_texture_kernel <<< grid, block >> > (gpu_dst, dstW, dstH, ratio);
    T.print("  核心計算");

    // 取出GPU值
    T.start();
    gpu_dst.memcpyOut(dst, dstSize);
    T.print("  GPU 取出資料");
}

總結

可以看出大部分時間其實是消耗再第一次new空間的時候，至於記憶體的複製貌似也沒消耗太多時間，建議圖可以直接讀進共享記憶體就好了，不需要先用主機記憶體讀然後再複製進去，從一開始的時候就宣告共享記憶體。

texture times = 0.327s
share times = 0.4s
CPU times = 2.578s

範例用的原始碼：https://github.com/hunandy14/Cuda_ImgBilinear

​