@作者: 機(jī)器學(xué)習(xí)算法 @迪吉老農(nóng)
代碼地址：https://github.com/yandili/forge_load

1. 背景需求

最近組內(nèi)的GPU利用率一直被警告，說是利用率過低。其實(shí)GPU這件事和CPU還是有區(qū)別的。

第一個問題是內(nèi)存限制。CPU的話，可以平行的跑很多程序，這樣利用率就上去了。但GPU很大程度上受限于內(nèi)存。如果內(nèi)存只能裝2個進(jìn)程，再想運(yùn)行更多的程序也沒有辦法。

第二個問題是，CPU一般可以通過復(fù)制進(jìn)程來提高利用率，每個進(jìn)程占用一個CPU核，就可以按任意的比例提高總體利用率。但是GPU的訓(xùn)練任務(wù)跑起來的時候，經(jīng)常一個程序就100%占用了。如果用這種方式占用空閑GPU，別的正常的程序就只能等待了。

不過既然上面要求了，我們也得做。就考慮兩個方面的要求，

• 占用盡可能小的內(nèi)存。
• 控制單進(jìn)程的GPU資源占用比例。

方案一（廢棄）

啟動一個接口程序，類似于圖像分類任務(wù)，模擬用戶請求，通過增加請求量的方式來增加負(fù)載。

缺點(diǎn)：

• 加載模型的話，會消耗一定比例的內(nèi)存，我印象比較小的模型也有好幾百M(fèi)B

方案二（采用）

研究一下NVIDIA提供的CUDA接口，直接調(diào)用CUDA編寫GPU程序，進(jìn)行簡單的并行計算來占用GPU核。這樣基本不消耗內(nèi)存，并且可以精確控制GPU核心數(shù)。

2. 調(diào)研

這一塊兒簡單了解一下CUDA和python的接口。撿了幾個主要概念看了一下。

CUDA的基本概念

CUDA的C語言文檔

• GPU的核心是Streaming Multiprocessors（sm），數(shù)量成千上萬。核心有三個概念

• a hierarchy of thread groups, shared memories, and barrier synchronization

• GPU和CPU

• hierachy of thread groups是說，計算任務(wù)都是按照矩陣的格式思考。sm的排列可以理解是矩陣，一個子矩陣叫g(shù)rid，grid的行叫block，具體的元素是thread?？偟挠嬎阗Y源占用就是從大的矩陣?yán)镆?guī)劃出來的子矩陣中的所有thread。這個比例基本就對應(yīng)著GPU的利用率。

• Grid

• shared memories是說，block內(nèi)部的幾個thread，在計算的時候是有一個內(nèi)部高速cache，如果好幾個thread要重復(fù)讀同一條數(shù)據(jù)，那最好在算法里把這幾個sm放到一個block里。這個我也沒仔細(xì)看，有個矩陣乘法運(yùn)算的例子，再補(bǔ)充。

• barrier synchronization是說，block內(nèi)部的幾個thread，是可以等待一起完成的？也沒仔細(xì)看。

• 規(guī)劃出來的一個grid，所有的thread是同時拿到一個函數(shù)，同時執(zhí)行。這個函數(shù)在CUDA語義下叫kernel。函數(shù)里面有變量可以方便每個thread定位自己所在的行數(shù)和列數(shù)。每個thread通過這個行數(shù)和列數(shù)，判斷自己需要執(zhí)行的操作。

• 之前有同事告訴我，GPU是一個進(jìn)程獨(dú)占全部thread。這個問題現(xiàn)在看來，可能也不一定對，還是看申請了多少thread，剩下的應(yīng)該不被占用。

• 數(shù)據(jù)是需要在CPU和GPU之間傳遞的，有兩份存在。

numba程序?qū)嶒?/h3>

python想要調(diào)用cuda的功能，需要借助numba。本質(zhì)上numba是通過預(yù)編譯python代碼加速矩陣運(yùn)算的。numba提供了一個cuda的接口。CUDA的python文檔。

比如，下面的python程序，將一個4*128的矩陣，并行填上數(shù)字。其中的grid, threadIdx都是可以定位用。

from numba import cuda
@cuda.jit
def my_kernel(io_array):
    # pos = cuda.grid(1) 是thread個數(shù)的一個index, 比如按照后面的配置，2*128=256個
    # tx = cuda.threadIdx.x 是每個block內(nèi)部0-128的index
    # assert pos == cuda.threadIdx.x + cuda.threadIdx.y * cuda.blockDim.x
    pos = cuda.grid(1)
    tx = cuda.threadIdx.x 
    if pos < io_array.size:
        io_array[pos] += tx # do the computation

blocks = 2
threadings = 128
data = np.zeros(512)
# 在運(yùn)行時指定，用多少thread執(zhí)行函數(shù)，其中的方括號的格式看起來比較奇怪，是對應(yīng)的在C語言接口里的
#  MatAdd<<<numBlocks, threadsPerBlock>>>(A, B, C); 這種
# <<<...>>>
my_kernel[blockspergrid, threadsperblock](data)

輸出結(jié)果如下，

[  0.   1.   2.   3.   4.   5.   6.   7.   8.   9.  10.  11.  12.  13.
  14.  15.  16.  17.  18.  19.  20.  21.  22.  23.  24.  25.  26.  27.
  28.  29.  30.  31.  32.  33.  34.  35.  36.  37.  38.  39.  40.  41.
  42.  43.  44.  45.  46.  47.  48.  49.  50.  51.  52.  53.  54.  55.
  56.  57.  58.  59.  60.  61.  62.  63.  64.  65.  66.  67.  68.  69.
  70.  71.  72.  73.  74.  75.  76.  77.  78.  79.  80.  81.  82.  83.
  84.  85.  86.  87.  88.  89.  90.  91.  92.  93.  94.  95.  96.  97.
  98.  99. 100. 101. 102. 103. 104. 105. 106. 107. 108. 109. 110. 111.
 112. 113. 114. 115. 116. 117. 118. 119. 120. 121. 122. 123. 124. 125.
 126. 127.   0.   1.   2.   3.   4.   5.   6.   7.   8.   9.  10.  11.
  12.  13.  14.  15.  16.  17.  18.  19.  20.  21.  22.  23.  24.  25.
  26.  27.  28.  29.  30.  31.  32.  33.  34.  35.  36.  37.  38.  39.
  40.  41.  42.  43.  44.  45.  46.  47.  48.  49.  50.  51.  52.  53.
  54.  55.  56.  57.  58.  59.  60.  61.  62.  63.  64.  65.  66.  67.
  68.  69.  70.  71.  72.  73.  74.  75.  76.  77.  78.  79.  80.  81.
  82.  83.  84.  85.  86.  87.  88.  89.  90.  91.  92.  93.  94.  95.
  96.  97.  98.  99. 100. 101. 102. 103. 104. 105. 106. 107. 108. 109.
 110. 111. 112. 113. 114. 115. 116. 117. 118. 119. 120. 121. 122. 123.
 124. 125. 126. 127.   0.   0.   0.   0.   0.   0.   0.   0.   0.   0.
   0.   0.   0.   0.   0.   0.   0.   0.   0.   0.   0.   0.   0.   0.
   0.   0.   0.   0.   0.   0.   0.   0.   0.   0.   0.   0.   0.   0.
   0.   0.   0.   0.   0.   0.   0.   0.   0.   0.   0.   0.   0.   0.
   0.   0.   0.   0.   0.   0.   0.   0.   0.   0.   0.   0.   0.   0.
   0.   0.   0.   0.   0.   0.   0.   0.   0.   0.   0.   0.   0.   0.
   0.   0.   0.   0.   0.   0.   0.   0.   0.   0.   0.   0.   0.   0.
   0.   0.   0.   0.   0.   0.   0.   0.   0.   0.   0.   0.   0.   0.
   0.   0.   0.   0.   0.   0.   0.   0.   0.   0.   0.   0.   0.   0.
   0.   0.   0.   0.   0.   0.   0.   0.   0.   0.   0.   0.   0.   0.
   0.   0.   0.   0.   0.   0.   0.   0.   0.   0.   0.   0.   0.   0.
   0.   0.   0.   0.   0.   0.   0.   0.   0.   0.   0.   0.   0.   0.
   0.   0.   0.   0.   0.   0.   0.   0.   0.   0.   0.   0.   0.   0.
   0.   0.   0.   0.   0.   0.   0.   0.   0.   0.   0.   0.   0.   0.
   0.   0.   0.   0.   0.   0.   0.   0.   0.   0.   0.   0.   0.   0.
   0.   0.   0.   0.   0.   0.   0.   0.   0.   0.   0.   0.   0.   0.
   0.   0.   0.   0.   0.   0.   0.   0.   0.   0.   0.   0.   0.   0.
   0.   0.   0.   0.   0.   0.   0.   0.   0.   0.   0.   0.   0.   0.
   0.   0.   0.   0.   0.   0.   0.   0.]

可以看到512維的數(shù)據(jù)只有前兩個block被更新了，可以理解邏輯是這樣的：

傳入一個待寫入的內(nèi)存空間，如果需要寫滿所有的列，需要len(blocks) \* len(threads per block)的thread，總數(shù)要超過data的行列數(shù)。

每個thread，會拿到自己的位置，并判斷自己是否執(zhí)行（比如是否在data范圍內(nèi)）。

3. 具體實(shí)現(xiàn)

實(shí)現(xiàn)GPU利用率的提升

這個完全可以通過block和thread per block的數(shù)量來控制。并且，我們并不需要開一個很大的內(nèi)存空間，最終thread的數(shù)量和data的大小不需要一致。一個thread的kernel即使不做任何事情，也是會被鎖定占用的。

控制GPU利用率在給定范圍

理論上，如果我們知道GPU總共可以提供多少thread數(shù)，我們按比例鎖定就好，不過沒有查到這個總資源量。考慮使用動態(tài)的方式調(diào)節(jié)這個數(shù)量。我們使用Worker來提升利用率，用Monitor來監(jiān)控利用率。程序啟動的時候，設(shè)定一個初始的thread數(shù)，比如1000。當(dāng)利用率不足的時候，按固定比例提升thread數(shù)；當(dāng)利用率低于預(yù)期時，按固定比例降低thread數(shù)。

避免占用正常程序資源

Monitor檢測到負(fù)載比較大之后，會自動休眠。只有沒有程序負(fù)載之后，才啟動。

下面是設(shè)定占用50%的程序日志，可以看到負(fù)載值load和thread數(shù)的增加比例multiplier的變化，

threadsperblock: 128, blockspergrid: 4
Monitor started: True
('Initial average load', 0.0)
Run for 10s with load 0.0 and multiplier 1000
Adjusted speed: boost
Run for 10s with load 12.4 and multiplier 1200.0
Adjusted speed: boost
Run for 10s with load 14.4 and multiplier 1440.0
Adjusted speed: boost
Run for 10s with load 15.0 and multiplier 1728.0
Adjusted speed: boost
Run for 10s with load 16.0 and multiplier 2073.6
Adjusted speed: boost
Run for 10s with load 17.6 and multiplier 2488.32
Adjusted speed: boost
Run for 10s with load 24.4 and multiplier 2985.984
Adjusted speed: boost
Idle for 5s with load 69.0
Run for 10s with load 4.8 and multiplier 3583.1808
Adjusted speed: boost
Run for 10s with load 28.8 and multiplier 4299.81696
Adjusted speed: boost
Run for 10s with load 33.6 and multiplier 5159.780352
Adjusted speed: boost
Run for 10s with load 39.6 and multiplier 6191.7364224
Idle for 5s with load 56.00000000000001
Run for 10s with load 0.0 and multiplier 6191.7364224
Run for 10s with load 47.2 and multiplier 6191.7364224
Run for 10s with load 45.6 and multiplier 6191.7364224
Adjusted speed: boost
Run for 10s with load 43.8 and multiplier 7430.08370688
Run for 10s with load 54.8 and multiplier 7430.08370688
Idle for 5s with load 56.00000000000001
Run for 10s with load 8.6 and multiplier 7430.08370688
Idle for 5s with load 56.00000000000001
Run for 10s with load 0.0 and multiplier 7430.08370688

最終負(fù)載會穩(wěn)定在50%左右，內(nèi)存占用為110MB。

image-20181107204631853.png

4. 部署

考慮到不同機(jī)器的配置不同，尤其是cuda和cudnn的版本不同，考慮還是通過docker方式封裝。

描述

按照給定比例占用GPU資源

• GPU內(nèi)存占用為120MB

• 占用比例按照使用情況動態(tài)調(diào)節(jié)，如果有訓(xùn)練程序，會自動讓給訓(xùn)練程序使用

• 默認(rèn)占用50%左右的GPU，閑置時間長的話，會增加到60%

安裝需求

機(jī)器上需要預(yù)先裝有cuda，并且下面的命令已經(jīng)可以輸出

nvidia-docker run -it nvidia/cuda:8.0-cudnn5-devel nvidia-smi

配置參數(shù)

NV_GPU=0  # 顯卡號：0或者1

啟動命令

NV_GPU={NV_GPU} nvidia-docker run -d\
    --log-driver=json-file --log-opt max-size=3m --log-opt max-file=3 \
    --name forge_load_gpu_{NV_GPU} \
    registry.api.weibo.com/forge_load/forge_load:0.1-gpu

終止程序

docker rm -f forge_load_gpu_${NV_GPU} 

版權(quán)聲明

以上文章為本人@迪吉老農(nóng)原創(chuàng)，首發(fā)于簡書，文責(zé)自負(fù)。文中如有引用他人內(nèi)容的部分（包括文字或圖片），均已明文指出，或做出明確的引用標(biāo)記。如需轉(zhuǎn)載，請聯(lián)系作者，并取得作者的明示同意。感謝。