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\ ):

OpenCV  現在有支援 CUDA 與 OpenCL, 但是官網預先編譯過的程式庫並未包含 CUDA, 所以要自行從原始檔進行編譯.
1. 先從 OpenCV官網下載原始檔, 然後解壓縮.

由於 CUDA 對 Visual Stidio 的 Integration 會檢查版本, 如果安裝 CUDA 之後更新升級 Visual Stdio, 可能導致因版本錯誤而出現以下訊息:
    #error -- unsupported Microsoft Visual Studio version! Only the versions 2012, 2013, 2015 and 2017 are supported!

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

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

這裡使用兩種方式:

• 使用 windows的 PerformanceCounter, 只要 include <Windows.h> 即可使用
• 使用 cudaEvent

arraySize受限於host與gpu的記憶體容量, 可自行調整測試, repeat為重複執行次數. 如果在 main() 中重複呼叫 addWithCuda(), 會因為記憶體搬移花費大量時間, 反而可能比CPU還慢, 因此直接在 addWithCuda() 中重複呼叫 kernel去計算. 資料陣列可宣告於 globalㄝ如果放在local 最好宣告 static, 否則  arraySize 太大會造成 stack overflow.

CudaInfo(): 列印Cuda相關能力

以下程式碼展示 host memory與device memory之間的關係, 變數命名 host 開頭表示是在 host memory, dev 開頭表示是在 device memory. 一開始初始化 host_a[], 然後用cudaMemcpy() 複製至 dev_a[], 平行呼叫 kernel 函數 shared() 後, 將 dev_a[] 複製至 x[], x是shared memory, 可在同一個 block 內的所有 thread 使用.

在 CUDA 中, 以軟體觀點來看, 一個 grid 由許多 block 組成, 一個 block 由許多 thread 組成, 每一個 kernel 函數執行時, 可以從 threadIdx 與 blockIdx 取得編號, 但是並沒有 gridIdx, 這表示每一個平行處理的執行, 是在同一個 grid 下進行. BlockSize 與 ThreadSize 的數量受限於 GPU 的硬體架構, 平行處理的函數呼叫語法是:

    函數名稱 < << BlockSize, ThreadSize>> > (函數參數)

最近開始研究CUDA，準備做演算法加速， 因此生出這篇筆記。因CUDA 是叫用 nvcc 進行編譯，CUDA函數無法被C#直接呼叫，在同一個CUDA專案中只能用C或C++來呼叫 CUDA，因此必須在CUDA專案中加一個C/C++的wrapper函數，將CUDA專案包裝成C語言的DLL檔，然後在C#中用DllImport呼叫C函式來轉給CUDA計算。以下實作以Visual Studio 2017 Community與 CUDA 9.2 SDK為例，完整程式碼在https://github.com/ghostyguo/CudaDotNet。