用OLED螢幕播放視訊(3): 使用cuda程式設計加速視訊處理

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

這是此係列文章的第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