因為影像數據都是存放在GPU上，我們要修改就不得不使用CUDA，但是如果我們先把數據複製到Host，修改完才複製回GPU，一來一回就大幅增加了延遲，所以直接使用CUDA Kernel修改數據是最快的途徑。CUDA Kernel就是在GPU上運行的函式，他的優勢就是可以有上千個執行緒並行處理數據，這個特性在處理像素上就有著巨大的優勢，因為大多數的情境都可以獨立修改每個像素。如果是初次接觸CUDA Kernel的人容易畏懼學習，但其實沒有大家想像的那麼困難，今天就帶大家來寫我們這系列文章最關鍵的修改顏色函式

科普一下CUDA編程

CUDA的成功有部分歸因於它讓調用GPU資源這件事情變得非常簡單，我們可以用C/C++的語法直接撰寫kernel，透過專屬的nvcc編譯器把程式編譯為.o檔案，這個.o檔案可以直接被C++的linker使用。所以我們可以跟平常一樣寫.h宣告C/C++的函式，在.cu內寫出kernel並且調用

手把手打造Kernel

寫程式碼

因為kernel需要另外的編譯器，為了避免跟原本的元件混在一起，所以另外寫一份.h和.cu。先寫color_magic.h

#include <stdint.h>
#include <stddef.h>
#include <cuda_runtime_api.h>

#ifdef __cplusplus
extern "C" {
#endif

void perform_color_magic_rgba(uint32_t width, uint32_t height, size_t pitch, uint8_t* data, cudaStream_t stream);

#ifdef __cplusplus
}
#endif

因為nvcc對C/C++的標準支援比較完善，所以建議避開GLib，而cuda_runtime_api.h是為了cudaStream_t引用

接著寫color_magic.cu

#include "color_magic.h"

__global__ void perform_color_magic_rgba_kernel(uint32_t width, uint32_t height, size_t pitch, uint8_t* data)
{
    uint32_t x = blockIdx.x * blockDim.x + threadIdx.x;
    uint32_t y = blockIdx.y * blockDim.y + threadIdx.y;

    if (x >= width || y >= height)
    {
        return;
    }

    uchar4* row = (uchar4*)(data + pitch * y);
    uchar4 pixel = row[x];

    /* Change color order from RGBA to GBRA */
    row[x] = make_uchar4(pixel.y, pixel.z, pixel.x, pixel.w);
}

bool perform_color_magic_rgba(uint32_t width, uint32_t height, size_t pitch, uint8_t* data, cudaStream_t stream)
{
    if (width == 0 ||
        height == 0 ||
        pitch < (size_t)width * 4 ||
        !data)
    {
        return false;
    }

    /* Hardcode block size for simple, you can optimize by yourself */
    dim3 blockDim{16, 16};
    dim3 gridDim{(width + blockDim.x - 1) / blockDim.x,
                 (height + blockDim.y - 1) / blockDim.y};

    perform_color_magic_rgba_kernel<<<gridDim, blockDim, 0, stream>>>(
        width, height, pitch, data);

    cudaError_t err = cudaGetLastError();
    if (err != cudaSuccess)
    {
        return false;
    }

    return true;
}

最後是修改Makefile，把nvcc加上去，並且讓它負責處理.cu的編譯工作

TARGET := libgstironmancolormagic.so
CUDA_DIR := /usr/local/cuda
DS_DIR := /opt/nvidia/deepstream/deepstream

CC := gcc
NVCC := $(CUDA_DIR)/bin/nvcc

PKGS := gstreamer-1.0 gstreamer-base-1.0

CFLAGS := -Wall -Wextra -fPIC $(shell pkg-config --cflags $(PKGS)) \
    -I$(CUDA_DIR)/include \
    -I$(DS_DIR)/sources/includes

# My GPU's compute capability is 8.6, you can modify for your GPU
NVCC_CFLAGS := -arch=sm_86 \
    -Xcompiler -Wall,-Wextra,-fPIC

LDFLAGS := -shared $(shell pkg-config --libs $(PKGS)) \
    -L$(CUDA_DIR)/lib64 -lcudart \
    -L$(DS_DIR)/lib -lnvds_meta -lnvdsgst_meta \
    -Wl,-rpath,"$(CUDA_DIR)/lib64:$(DS_DIR)/lib"

SRCS := gstironmancolormagic.c
CU_SRCS := color_magic.cu
OBJS := $(SRCS:.c=.o) $(CU_SRCS:.cu=.o)

all: $(TARGET)

$(TARGET): $(OBJS)
    $(CC) -o $@ $^ $(LDFLAGS)

%.o: %.c
    $(CC) $(CFLAGS) -c $< -o $@

%.o: %.cu
    $(NVCC) $(NVCC_CFLAGS) -c $< -o $@

clean:
    rm -f $(OBJS) $(TARGET)

.PHONY: all clean

程式解說

輸入檢查

    if (width == 0 ||
        height == 0 ||
        pitch < (size_t)width * 4 ||
        !data)
    {
        return false;
    }

首先perform_color_magic_rgba要負責檢查輸入，確保輸入安全才傳入kernel，因為kernel是非同步啟動的，要把kernel的錯誤傳回host是很麻煩的

啟動kernel

    /* Hardcode block size for simple, you can optimize by yourself */
    dim3 blockDim{16, 16};
    dim3 gridDim{(width + blockDim.x - 1) / blockDim.x,
                 (height + blockDim.y - 1) / blockDim.y};

啟動kernel前我們要分配好需要調度的資源，我這邊blockDim為了方便示範是寫死的，實際上這個會受到硬體可用的執行緒數量影響。gridDim則是2D影像分配的標準公式，照抄就可以了。調度資源的分配屬於CUDA的範圍，建議大家還是花時間去閱讀官方的文件，我就不獻醜了(畢竟我自己也說不清楚)

perform_color_magic_rgba_kernel<<<gridDim, blockDim, 0, stream>>>(
        width, height, pitch, data);

再來我們想要支援非同步操作的話，啟動時就要帶入4個參數。第三個參數是sharedMemBytes但我們沒有需要共享資源，所以填入0。第4個參數就是stream，這樣就可以實現非同步運行kernel

kernel如何運作

我們看perform_color_magic_rgba_kernel

    uint32_t x = blockIdx.x * blockDim.x + threadIdx.x;
    uint32_t y = blockIdx.y * blockDim.y + threadIdx.y;

這是標準的2D處理取得目前執行緒位置的方法，舉個例子，如果處理一張2x2的影像啟動了4個執行緒，那麼就會有(0,0)、(0,1)、(1,0)、(1,1)等4個執行緒的位置，上面的x、y也是同理，我們就可以拿這個位置去找出這個執行緒應該處理的像素是哪一個

    if (x >= width || y >= height)
    {
        return;
    }

實際啟動的執行續數量可能會因為效率考量而超過影像的尺寸，所以這邊要進行邊界安全的檢查

    uchar4* row = (uchar4*)(data + pitch * y);
    uchar4 pixel = row[x];

    /* Change color order from RGBA to GBRA */
    row[x] = make_uchar4(pixel.y, pixel.z, pixel.x, pixel.w);

最後就是把像素的顏色數值重新排序一下，uchar4是CUDA定義好的結構，使用這個的原因是在CUDA當中，透過向量化指令操作 4 bytes 的資料只需 1 次指令(如 LDG.E.32)；而如果分開讀取 4 個單一 byte (uchar) 的資料，則需要 4 次指令(如 LDG.E.8)。總之，用uchar4就是比4個uint8_t分開操作還要快

檢查啟動的錯誤

    cudaError_t err = cudaGetLastError();
    if (err != cudaSuccess)
    {
        return false;
    }

離開函式前記得要檢查啟動kernel是否有成功，如果失敗的話cudaGetLastError就會回傳非cudaSuccess錯誤碼

導入transform_ip

新增引用

/* some code */

#include <gst/gst.h>
#include <gst/base/gstbasetransform.h>
#include <gst/video/video.h>
#include <gstnvdsmeta.h>
#include <nvbufsurface.h>

#include "gstironmancolormagic.h"
#include "color_magic.h"

到gstironmancolormagic.c加上color_magic.h的引用

呼叫perform_color_magic

到transform_ip內修改這一段

    for (NvDsFrameMetaList* l_frame = batch_meta->frame_meta_list; l_frame; l_frame = l_frame->next)
    {
        NvDsFrameMeta* frame_meta = l_frame->data;
        /* Remeber to use frame_meta->batch_id to get the image data of the frame */
        NvBufSurfaceParams* surf_params = &surf->surfaceList[frame_meta->batch_id];

        if (surf_params->colorFormat != NVBUF_COLOR_FORMAT_RGBA)
        {
            GST_ELEMENT_ERROR(self, RESOURCE, SETTINGS,
                ("NvBufSurfaceParams is not RGBA format"),
                ("'surf_params->colorFormat' is %" G_GINT32_FORMAT,
                (gint)surf_params->colorFormat));
            ret = GST_FLOW_ERROR;
            goto unmap_buffer;
        }

        /* Do color magic on `surf_params` asynchronously */
        if (!perform_color_magic_rgba(surf_params->width, surf_params->height,
                surf_params->pitch, surf_params->dataPtr, self->stream))
        {
            GST_ELEMENT_ERROR(self, LIBRARY, FAILED,
                ("Failed to perform color magic"),
                (NULL));
            ret = GST_FLOW_ERROR;
            goto sync_tasks;
        }
    }

輸入的顏色必須要是RGBA，所以先檢查surf_params->colorFormat，再來就是呼叫perform_color_magic_rgba，整個流程非常簡單