用OLED屏幕播放視頻(3): 使用cuda編程加速視頻處理

下麵的系列文章記錄瞭如何使用一塊linux開發扳和一塊OLED屏幕實現視頻的播放:

1. 項目介紹
2. 為OLED屏幕開發I2C驅動
3. 使用cuda編程加速視頻處理

這是此系列文章的第3篇, 主要總結和記錄瞭如何使用cuda編程釋放GPU的算力. 在此之前嘗試過使用python調用opencv直接處理視頻數據, 但使用之後發現處理過程效率不高, 處理時間偏長. 後來想到還有一塊顯卡沒利用起來, 畢竟在前司見證了某國產GPGPU晶元從立項, 到流片再到回片驗證的整個過程, cuda編程也算是傳統藝能了. 最終效果看下麵的視頻:

跳轉到6:48, 直接觀看演示

1). 要用GPU做什麼

這裡不會介紹cuda的編程模型, cuda開發工具的使用等, 這部分內容可以參考cuda的官方文檔, 學習cuda編程的話, 看這個文檔就足夠了.

原始的視頻文件, 每幀畫面的解析度一般不會和我們的屏幕尺寸128x64匹配, 並且視頻是彩色的, 使用的OLED屏幕只能顯示黑白圖像. 所以視頻的數據必須經過resize和灰度處理之後才能發送給beaglebone black板子連接的OLED屏幕, 這部分視頻處理工作就是在GPU上進行的.

在host machine上的python程式使用opencv讀取視頻文件中的每一幀, 通過socket發送給cuda程式; cuda程式處理完數據之後, 再通過socket把數據發送給beagle board上的用戶態程式; beagle board上的用戶態程式, 把一幀數據寫入屏幕, 完成繪製.

2). kernel函數的演算法實現

下麵是kernel函數的部分代碼, oframe, ow, oh, 分別表示原始畫面數據, 原始的寬度和高度, nframe, nw, nh分別表示處理之後的畫面數據, 新的寬度和高度.

kernel中的resize操作, 使用最近臨方式, (i, j)是新畫面中的像素位置, 計算得到對應的原始畫面像素位置(oi, oj), 取出原始的rgb值, 使用公式計算出亮度, 最後根據閾值確定(i, j)這個像素的亮滅.

__global__ void resize_frame_kernel(unsigned char *oframe, int ow, int oh,
				    unsigned char *nframe, int nw, int nh,
				    int threshold, unsigned int *locks)
{
	for (int i = blockDim.x * blockIdx.x + threadIdx.x; i < nw;
	     i += blockDim.x * gridDim.x) {
		for (int j = blockDim.y * blockIdx.y + threadIdx.y; j < nh;
		     j += blockDim.y * gridDim.y) {
			int oi = i * ow / nw;
			int oj = j * oh / nh;

			unsigned char b = oframe[oj * ow * 3 + oi * 3];
			unsigned char g = oframe[oj * ow * 3 + oi * 3 + 1];
			unsigned char r = oframe[oj * ow * 3 + oi * 3 + 2];

			unsigned char brightness =
				r * 0.3 + g * 0.59 + b * 0.11;
			brightness = brightness >= threshold ? 1 : 0;
			brightness = brightness << (j % 8);
			// 以下代碼實現了一個自旋鎖
			bool leaveloop = false;
			while (!leaveloop) {
				if (atomicExch(&locks[j / 8 * nw + i], 1u) ==
				    0u) {
					nframe[j / 8 * nw + i] |= brightness;
					leaveloop = true;
					atomicExch(&locks[j / 8 * nw + i], 0u);
				}
			}
		}
	}
}

3). kernel函數中的併發問題

在上面的代碼清單中使用原子交換指令atomicExch實現了一個自旋鎖. 在kernel函數中使用鎖是因為, nframe的大小是128x8位元組, 屏幕解析度是128x64, nframe的每個bit控制一個像素, 當kernel中更新nframe時, 可能同時有多個線程想更新nframe中的同一個位元組. 關於這個自選鎖中while迴圈的寫法, 可以參考stack overflow.

4). 文末推廣

歡迎關註我的B站賬號, 或者加入QQ群838923389, 一起研究電腦底層技術, 一起搞事情:P