No More Codes

為了瞭解 nVidia 關於 JPEG 的支援, 蒐集一些資料, 做個紀錄

程式碼來源: https://github.com/NVIDIA/cuda-samples

ffmpeg 使用 cuda 編碼器的參數範例:
ffmpeg -i "concat:VTS_01_1.VOB|VTS_01_2.VOB|VTS_01_3.VOB|VTS_01_4.VOB" -c:v h264_nvenc -preset slow -crf 22 -c:a copy "output.mp4"

使用 cudaFilter 的範例 (須 include cudafilters.hpp)

OpenCV 的 GpuMat 類別, 可以在 device 的 global memory 建立 Mat 物件, 可以被 kernel 函數直接存取, 以下以 sobel 運算子為例, 將 srcMat (位於host memory) 載入影像後, 轉換至 srcGMat (位於device memory), 經處理後放至 destGMat (位於device memory), 然後下載回 destMat (位於host memory)再顯示出來.
程式中使用 OpenCV 定義於 cuda_types.hpp 的 PtrStep 的樣板結構, 代入定義於 WIndows.h 的 byte, 這樣比較方便. 如果要用 cuda 編譯器 nvcc 內建的  uchar1 , 查看 vector_types.h, 它是一個結構 (向量):

取得 CUDA 版的 OpenCV 後, 使用 OpenCV 前, 先建立環境變數 OPENCV_DIR, 使其指向 opencv 目錄 ( C:\opencv\build\ ):
建立名稱為 OpenCvExample 的 CUDA 專案, CUDA精靈會自動建立一個 kernel.cu 檔 :

OpenCV  現在有支援 CUDA 與 OpenCL, 但是官網預先編譯過的程式庫並未包含 CUDA, 所以要自行從原始檔進行編譯.
1. 先從 OpenCV官網下載原始檔, 然後解壓縮.
2. 下載 CMake 安裝檔, 並完成安裝.
啟動 CMake 後, 在上方先輸入解壓縮後的目錄以及輸出目錄 (一般習慣為專案目錄下的build),並勾選 Grouped 與 Advanced, 然後按 Configure 按鈕, 就會讀取專案的設定檔, 在還未 Generate 前, 選項是紅色的:

由於 CUDA 對 Visual Stidio 的 Integration 會檢查版本, 如果安裝 CUDA 之後更新升級 Visual Stdio, 可能導致因版本錯誤而出現以下訊息:
    #error -- unsupported Microsoft Visual Studio version! Only the versions 2012, 2013, 2015 and 2017 are supported!
這個問題可以透過修改 C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v9.2\include\crt\host_config.h 的內容來解決, 只要將檢查版本的敘述修正一下就可以 :

影像處理 histogram 等化處理時, 把 histogram 視為機率分布函數 (pdf), 經常將其積分為累積分布函數 cdf, 例如:

考慮一個典型的總和運算 :

     int sum = 0;

     for (i = 1; i < N; i++) {

         sum += i;

   }

如果 kernel 執行時, 不同的 thread 對同一個變數進行讀取與寫入, 會發生甚麼事呢? 當不同thread被同時執行時, 同時去讀取同一個變數, 得到相同的值, 如果把此值運算後同時又寫回去, 最後的值就不一定是多少了, 以下程式開啟 BlockSize*ThreadSize 個 thread 去執行把 global memory 陣列變數加1的動作 :

#include "cuda_runtime.h"

#include "device_launch_parameters.h"

#include <stdio.h>

 

#define ThreadSize 1000

#define BlockSize  10000

#define ArraySize 10

 

__global__ void incKernel(int *a)

{

     int i = (blockIdx.x*blockDim.x + threadIdx.x) % ArraySize;

     a[i] = a[i] + 1;

}

 

int main()

{

     int host_a[ArraySize];

     int *dev_a = 0;

     float elapsedTime;

 

     // setup performance meter from CUDA ----------

     cudaEvent_t start, stop;

     cudaEventCreate(&start);

     cudaEventCreate(&stop);

 

     cudaSetDevice(0);

     cudaMalloc((void**)&dev_a, ArraySize * sizeof(int));

 

     for (int run = 0; run < 10; run++) {

 

         cudaMemset(dev_a, 0, ArraySize * sizeof(int));         //clear

 

         cudaEventRecord(start, 0); //keep start time

         incKernel << <BlockSize, ThreadSize >> > (dev_a); //calculate

         cudaEventRecord(stop, 0); //keep stop time

         cudaEventSynchronize(stop); //wait stop event    

         cudaEventElapsedTime(&elapsedTime, start, stop); 

 

         cudaMemcpy(host_a, dev_a, ArraySize * sizeof(int), cudaMemcpyDeviceToHost);

         //Print result

         printf("run {%d}: ",run);

         for (int i = 0; i < ArraySize; i++) {

              printf("%d ", host_a[i]);

         }

         printf(" t=%f\n",elapsedTime);

     }

     //cudaDeviceSynchronize();

     getchar();

 

     cudaFree(dev_a);

     return 0;

 

}