基于CUDA的KNN算法并行化研究

劉端陽，鄭江帆，劉 志

(浙江工業(yè)大學(xué) 計(jì)算機(jī)科學(xué)與技術(shù)學(xué)院，杭州310023)

1 引 言

分類技術(shù)在許多領(lǐng)域得到廣泛應(yīng)用，比如人工智能，網(wǎng)絡(luò)入侵檢測等.分類技術(shù)是把未知分類的樣本數(shù)據(jù)經(jīng)過一系列模型之后分到若干已知分類的過程[1].它先通過已有的訓(xùn)練數(shù)據(jù)集進(jìn)行學(xué)習(xí)和分析，從這些數(shù)據(jù)的分布中找出分類的規(guī)律，逐步優(yōu)化分類的模型，從而判斷新數(shù)據(jù)屬于哪個(gè)類別.K最近鄰(KNN：K-Nearest Neighbor)算法[2]是最為流行的分類算法之一，廣泛應(yīng)用于數(shù)據(jù)分析、圖像處理、文本分類[3]等領(lǐng)域.KNN算法在距離計(jì)算階段需要計(jì)算每個(gè)測試數(shù)據(jù)點(diǎn)到所有訓(xùn)練數(shù)據(jù)點(diǎn)的距離，距離排序階段需要對距離向量進(jìn)行排序.如果任務(wù)的測試數(shù)據(jù)總量為m，訓(xùn)練數(shù)據(jù)總量為n，數(shù)據(jù)的維度為d，則距離計(jì)算的時(shí)間復(fù)雜度為O(mnd)，排序階段的時(shí)間復(fù)雜度為O(mnlogn).隨著測試數(shù)據(jù)集和訓(xùn)練數(shù)據(jù)集規(guī)模的增加，算法的時(shí)間成本也會成倍增長，無法滿足大數(shù)據(jù)時(shí)代下的快速計(jì)算需求.因此對于計(jì)算密集型的KNN算法，運(yùn)用CUDA(Compute Unified Device Architecture)并行化該算法，能夠節(jié)省大量計(jì)算時(shí)間.CUDA是NV-IDIA公司推出的通用并行計(jì)算架構(gòu)，它采用GPU(Graphics Processing Unit)來解決復(fù)雜計(jì)算問題.CUDA提供了硬件的直接訪問接口，并不依賴傳統(tǒng)的圖形API，實(shí)現(xiàn)GPU的訪問，降低了圖形處理器的編程難度.

國內(nèi)外已經(jīng)有很多基于CUDA架構(gòu)的KNN以及其他算法的并行化研究，Jian等[4]將輸入的數(shù)據(jù)分組，在每個(gè)組內(nèi)進(jìn)行并行計(jì)算，然后得到每個(gè)組的top-K元素隊(duì)列，最后合并每個(gè)組得到全局的K個(gè)近鄰，該方法降低了距離排序的時(shí)間復(fù)雜度，但是在距離的計(jì)算階段并行度不高.Garcia等[5,6]將KNN的暴力搜索算法并行移植到GPU中執(zhí)行，在距離計(jì)算階段使用了分塊的矩陣乘法，具有良好并行性，但是在距離排序階段時(shí)間復(fù)雜度較高.Dashti等[7]在面對高維度數(shù)據(jù)時(shí)合理使用線程加速距離計(jì)算，并且在GPU中并行實(shí)現(xiàn)了三種距離排序方法，但是當(dāng)訓(xùn)練集數(shù)據(jù)量很大時(shí)排序階段的加速效果有限.Masek等[8]和王裕民等[9]使用多塊GPU設(shè)備搭建了一個(gè)GPU集群，利用GPU的集群優(yōu)勢對KNN和卷積神經(jīng)網(wǎng)絡(luò)算法進(jìn)行加速，該方法可以處理更加龐大的數(shù)據(jù)集，不過提高了硬件成本.Liang等[10]提出cuKNN算法，采用數(shù)據(jù)流的方式對輸入數(shù)據(jù)進(jìn)行計(jì)算，利用共享內(nèi)存加速每個(gè)測試數(shù)據(jù)距離向量的排序，能夠批量處理輸入數(shù)據(jù)的同時(shí)也提升了性能.Barrientos等[11]提出基于CUDA的窮舉排序算法解決KNN問題，在K值較高的情況下也有良好的加速比.Mayekar等[12]提出一種并行的KNN方法快速解決字符識別的問題，并且能保持較高的準(zhǔn)確率.

基于上述國內(nèi)外的研究現(xiàn)狀，本文提出了新的KNN并行化算法(即GS_KNN).新算法充分利用了CUDA函數(shù)庫以及內(nèi)存體系，將算法的所有計(jì)算過程都移植到GPU中，發(fā)揮其良好的運(yùn)算能力.在運(yùn)算量最大的距離計(jì)算階段，將測試數(shù)據(jù)集和訓(xùn)練數(shù)據(jù)集存儲在CUDA全局內(nèi)存中，使用Cublas函數(shù)庫的通用矩陣乘計(jì)算，提高了運(yùn)算速度.在距離排序階段，針對k值的大小提出了兩種優(yōu)化策略，分別是基于k次最小值查找的最近鄰選擇和基于雙調(diào)排序的最近鄰選擇，降低了時(shí)間復(fù)雜度.在決定分類標(biāo)號階段，采用了CUDA內(nèi)部的原子加法操作，GPU全程參與計(jì)算，從而提高了整體性能.最后通過實(shí)驗(yàn)證明，在保證算法魯棒性的前提下，獲得了良好的加速效果.

2 背景知識

2.1 KNN算法

KNN算法是目前最為流行的分類算法之一，該算法是一種基于距離的分類，通過計(jì)算測試數(shù)據(jù)與訓(xùn)練數(shù)據(jù)之間的距離或者相似度來進(jìn)行分類，同一類別的內(nèi)部數(shù)據(jù)之間具有較高的相似度，而不同類別之間數(shù)據(jù)相似度較低.現(xiàn)有R={r1，r2，…，rn}是一個(gè)擁有n個(gè)點(diǎn)的輸入訓(xùn)練集，其中每個(gè)點(diǎn)的數(shù)據(jù)維度為d，即rj=(rj1，rj2，…，rjd)，而S=(s1，s2，…，sm)是與R在同一個(gè)維度空間上的輸入測試集.KNN算法的目標(biāo)是對于每一個(gè)si∈S，搜索距離si最近的k個(gè)R數(shù)據(jù)點(diǎn)，然后根據(jù)這k個(gè)數(shù)據(jù)點(diǎn)的類別判斷si的類別.如圖1所示，表示一個(gè)R大小為15，k等于4的查詢例子.

圖1 KNN查詢示例Fig.1 KNN query example

求解KNN算法最經(jīng)典的方法就是通過暴力搜索解決，具體算法過程如算法1所示.

算法1.基于CPU的串行KNN算法

輸入：包含n個(gè)點(diǎn)的帶有分類標(biāo)號的訓(xùn)練數(shù)據(jù)集R；包含m個(gè)點(diǎn)的無分類標(biāo)號的測試數(shù)據(jù)集S；最近鄰個(gè)數(shù)k.

輸出：m個(gè)測試數(shù)據(jù)的分類.

1)計(jì)算測試集S中點(diǎn)si到所有rj的距離.

2)對步驟1得到的n個(gè)距離進(jìn)行升序排序.

3)根據(jù)排序結(jié)果，選擇前k個(gè)最小距離，得到距離點(diǎn)si最近的k個(gè)訓(xùn)練數(shù)據(jù).

4)統(tǒng)計(jì)這k個(gè)訓(xùn)練數(shù)據(jù)的分類標(biāo)號，出現(xiàn)次數(shù)最多的那一個(gè)標(biāo)號即為點(diǎn)si的分類.

5)循環(huán)步驟1至4，一直到m個(gè)測試數(shù)據(jù)的分類計(jì)算完畢.

由于數(shù)據(jù)維度為d，步驟1(即距離計(jì)算階段)的時(shí)間復(fù)雜度為O(nmd)，步驟2(即排序階段)的時(shí)間復(fù)雜度為O(mnlogn).步驟4(即決定分類標(biāo)號階段)的時(shí)間復(fù)雜度為O(mk)，此值較小，可忽略不計(jì).當(dāng)訓(xùn)練數(shù)據(jù)集n很大時(shí)或者面對高維度數(shù)據(jù)時(shí)，運(yùn)用CUDA計(jì)算架構(gòu)并行化KNN算法通常可以取得很好的加速效率.

2.2 CUDA編程模型

CUDA是NVIDIA公司推出的一款通用并行計(jì)算架構(gòu)，CUDA架構(gòu)改變了傳統(tǒng)的GPU編程方式，它不需要將程序任務(wù)轉(zhuǎn)換為GPU圖形處理任務(wù)，降低了開發(fā)難度.CUDA架構(gòu)主要包括計(jì)算核心和存儲器體系.每個(gè)GPU計(jì)算核心由多個(gè)流多處理器 (stream multiprocessor，SM) 組成，每個(gè)SM由多個(gè)標(biāo)量處理器 (stream processor，SP) 組成.同時(shí)，GPU提供了可供線程訪問的多級存儲器.每個(gè)線程具有私有的本地存儲器，每個(gè)線程塊則有共享存儲器，它對塊內(nèi)所有線程共享數(shù)據(jù)，并且它的生命周期與塊相同，讀寫速度快.常量存儲器和紋理存儲器是只讀存儲器，可以針對不同的用途進(jìn)行優(yōu)化，全局存儲器則用于接收從CPU端傳入的數(shù)據(jù).

在CUDA編程模型中，GPU被認(rèn)為是能夠并行執(zhí)行大量線程的協(xié)處理器.一個(gè)簡單的源程序包括運(yùn)行在CPU上的主機(jī)端代碼和運(yùn)行在GPU上的內(nèi)核(kernel)代碼.Kernel函數(shù)通過__global__類型限定符定義，并通過Kernel<<>>(parameters)調(diào)用，blocks代表線程塊數(shù)量，threads代表每個(gè)線程塊中包含的線程數(shù)，parameters代表函數(shù)調(diào)用時(shí)的參數(shù)列表.計(jì)算密集型和數(shù)據(jù)并行的內(nèi)核函數(shù)代碼運(yùn)行在GPU上.所有線程被組成許多線程塊，每個(gè)線程塊中的線程會在同一個(gè)SM中并行執(zhí)行.同一個(gè)線程塊的線程可以通過塊內(nèi)的共享內(nèi)存共享數(shù)據(jù)，并且可以通過__syncthreads()函數(shù)同步數(shù)據(jù).在CUDA編程模型中另一個(gè)重要的概念就是線程束，它由32個(gè)并行的線程組成并且是每個(gè)SM處理器的基本調(diào)度單元.當(dāng)一個(gè)線程束執(zhí)行結(jié)束之后，SM處理器會接著調(diào)度另一個(gè)線程束繼續(xù)執(zhí)行，線程束一次執(zhí)行同一條指令，所以當(dāng)線程束中的32條線程有相同的執(zhí)行路徑時(shí)可以達(dá)到最大的加速效率.當(dāng)遇到條件分支時(shí)線程束中的線程產(chǎn)生不同的執(zhí)行路徑，這將會增加該線程束執(zhí)行指令的總時(shí)間.為了讓資源不浪費(fèi)，即warp內(nèi)的線程被充分利用，線程的數(shù)量經(jīng)常設(shè)置為32的整數(shù)倍.只有合理地分配線程塊以及線程的數(shù)量，才可以使GPU的利用率最大化.

3 基于CUDA的KNN算法并行優(yōu)化

分析串行KNN算法的流程可知，算法的時(shí)間開銷大部分都在距離計(jì)算階段和排序階段.因此本文著重在這兩個(gè)步驟上進(jìn)行并行優(yōu)化，并且在GP-U中進(jìn)行所有計(jì)算.以下四個(gè)小節(jié)詳細(xì)介紹了GS_KNN算法的優(yōu)化過程.

3.1 基于Cublas的距離矩陣計(jì)算

Cublas Library是基于NVIDIA的通用函數(shù)庫，它實(shí)現(xiàn)了Blas(基本線性代數(shù)子程序)，它允許用戶訪問GPU中的計(jì)算單元，在使用時(shí)在GPU中分配所需的矩陣向量空間并填充數(shù)據(jù)，然后調(diào)用其中的函數(shù)可以加速向量和矩陣等線性運(yùn)算.對于KNN算法中龐大的距離計(jì)算量，本文利用Cublas庫中的矩陣乘法函數(shù)cublasSgemm()來加速距離計(jì)算，該函數(shù)能實(shí)現(xiàn)向量與矩陣，矩陣與矩陣之間的運(yùn)算，并且擁有很好的加速比.

cublasSgemm()函數(shù)調(diào)用時(shí)有許多參數(shù)，調(diào)用方法為cublasSgemm(handle，transa，transb，m，n，k，*alpha，*A，lda，*B，ldb，*beta，*C，ldc)，在Cublas里面所有的矩陣都是按照列優(yōu)先方式存儲，lda，ldb，ldc代表矩陣A，B，C的行數(shù)，矩陣A的維度為m*k，B的維度為k*n，C的維度為m*n，transa和transb代表矩陣A和B是否轉(zhuǎn)置，函數(shù)實(shí)現(xiàn)的矩陣運(yùn)算為：

C=alpha*OP(A)*OP(B)+beta*C

(1)

為了滿足公式(1)，GS_KNN算法的距離度量選擇歐幾里得距離的平方.這樣，GS_KNN算法的距離計(jì)算可以轉(zhuǎn)化為：

|x-y|2=x2+y2-2x*y=alpha(x*y)+beta(x2+y2)

(2)

其中，alpha取值為-2，beta取值為1.在設(shè)計(jì)CUDA核函數(shù)時(shí)，如果有m個(gè)測試數(shù)據(jù)點(diǎn)和n個(gè)訓(xùn)練數(shù)據(jù)點(diǎn)，數(shù)據(jù)維度為d，則矩陣A為x，代表m*d維的測試數(shù)據(jù)點(diǎn)矩陣，矩陣B為y，代表n*d維的訓(xùn)練數(shù)據(jù)點(diǎn)矩陣，由于x和y都是按照列優(yōu)先方式存儲，實(shí)際上在函數(shù)運(yùn)算時(shí)矩陣A的維度為d*m，矩陣B的維度為d*n，為了得到m*n維的距離結(jié)果矩陣，需要將矩陣A進(jìn)行轉(zhuǎn)置，矩陣B不需要轉(zhuǎn)置，所以將transa這個(gè)參數(shù)的值設(shè)為CUBLAS_OP_T代表轉(zhuǎn)置，transb這個(gè)參數(shù)的值設(shè)為CUBLAS_OP_N代表不需轉(zhuǎn)置.由于矩陣A的行數(shù)為m，矩陣B的列數(shù)為n，矩陣A的列數(shù)為d，所以分別將函數(shù)參數(shù)列表中的m賦值為m，n賦值為n，k賦值為d.在調(diào)用函數(shù)之前先要計(jì)算矩陣C也就是x2+y2的值，對于x2算法開啟m個(gè)線程，若矩陣A在CUDA內(nèi)存中的地址為p，線程編號為threadId，每個(gè)線程計(jì)算地址在p+d*threadId至p+d*(threadId+1)上的數(shù)據(jù)平方并求和，即一個(gè)測試數(shù)據(jù)點(diǎn)在d個(gè)維度上的平方和.對于y2開啟n個(gè)線程，若矩陣B在CUDA內(nèi)存中的地址為q，線程編號為threadId，每個(gè)線程計(jì)算地址在q+d*threadId至q+d*(threadId+1)上的數(shù)據(jù)平方并求和，即一個(gè)訓(xùn)練數(shù)據(jù)點(diǎn)在d個(gè)維度上的平方和.得到以上兩個(gè)結(jié)果之后，在CUDA中繼續(xù)開啟n個(gè)線程，每個(gè)線程將一個(gè)訓(xùn)練數(shù)據(jù)點(diǎn)的平方和分別與m個(gè)測試數(shù)據(jù)點(diǎn)的平方和相加從而得到維度為m的向量，則n個(gè)線程計(jì)算得到m*n維的矩陣，即表達(dá)式為x2+y2的輸入的矩陣C.最后將矩陣A，B和C在CUDA全局內(nèi)存中的地址指針傳入cublasSgemm()函數(shù)，函數(shù)返回m*n維的距離結(jié)果矩陣，每一行向量代表一個(gè)測試數(shù)據(jù)點(diǎn)分別與n個(gè)訓(xùn)練數(shù)據(jù)點(diǎn)的距離.

3.2 基于k次最小值查找的最近鄰選擇

通過計(jì)算獲得每個(gè)測試數(shù)據(jù)點(diǎn)的距離向量后，需要進(jìn)行排序.雖然有不少研究把排序算法移植到CUDA中執(zhí)行，如Sintorn等[13]和Satish等[14]的研究，但算法的時(shí)間復(fù)雜度為O(nlogn)，并行度不高.通過分析KNN算法可知，參數(shù)k往往遠(yuǎn)小于訓(xùn)練集數(shù)據(jù)量n，所以對于每一個(gè)測試數(shù)據(jù)來說，沒有必要將它到訓(xùn)練集的所有距離都進(jìn)行排序，只需要獲得前k個(gè)最小距離即可.冒泡排序算法和選擇排序算法，每次排序都可以選出一個(gè)最小值，k次排序就可以獲得k個(gè)最小距離.雖然這兩種算法的時(shí)間復(fù)雜度只有O(kn)，但是它們卻無法并行執(zhí)行，因?yàn)槊看闻判蚨家蕾嚽懊嬖氐呐判蚪Y(jié)果.因此，冒泡排序算法和選擇排序算法都不適合CUDA平臺.

本節(jié)提出了一個(gè)基于k次最小值查找的最近鄰選擇方法，適合GPU的多線程環(huán)境，并且是一個(gè)基于比較網(wǎng)絡(luò)的模型，可以在同一時(shí)刻執(zhí)行多個(gè)不相關(guān)的比較操作.算法定義的一個(gè)比較器如圖2所示，比較器的輸入為兩個(gè)任意未經(jīng)過排序的數(shù)據(jù)，經(jīng)過比較器之后兩個(gè)輸出被排序，并且上端線路的元素大于或者等于下端線路的元素，一個(gè)比較器所執(zhí)行的時(shí)間為一個(gè)單位時(shí)間，時(shí)間復(fù)雜度為O(1).

圖2 K次最小值查找的比較器
Fig.2 Comparator of K-min search

對于一個(gè)測試數(shù)據(jù)，經(jīng)過距離計(jì)算之后得到它對n個(gè)訓(xùn)練數(shù)據(jù)的距離向量{d0，d1，…dn-2，dn-1}，需要從該距離向量中選擇出k(k≤n)個(gè)最小距離，根據(jù)排序網(wǎng)絡(luò)的比較器需要兩個(gè)輸入，算法將該向量以dn/2處的位置一分為二，然后d0和d0+?n/2」(符號?」表示向下取整)輸入比較器做判斷之后得到max{d0，d0+?n/2」}和min{d0，d0+?n/2」}.與此同時(shí)，d1和d1+?n/2」做比較得到max{d1，d1+?n/2」}和min{d1，d1+?n/2」}，dn/2之前的一個(gè)元素d?n/2」-1和d?n/2」-1+?n/2」比較得到max{d?n/2」-1，d?n/2」-1+?n/2」}和min{d?n/2」-1，d?n/2」-1+?n/2」}，當(dāng)n為奇數(shù)時(shí)，最后一個(gè)數(shù)dn-1無須經(jīng)過比較自動進(jìn)入下一輪迭代.這些兩兩元素之間的比較是互不相關(guān)的，符合CUDA并行執(zhí)行的特性.在經(jīng)過第一輪比較器之后會產(chǎn)生兩個(gè)子向量l1和l2，如果n為偶數(shù)，則向量如下所示：

l1={max{d0，d0+?n/2」}，max{d1，d1+?n/2」}，…
max{d?n/2」-1，d?n/2」-1+?n/2」}}

(3)

l2={min{d0，d0+?n/2」}，min{d1，d1+?n/2」}，…
min{d?n/2」-1，d?n/2」-1+?n/2」}}

(4)

如果n為奇數(shù)，則子向量l2為：

l2={min{d0，d0+?n/2」}，min{d1，d1+?n/2」}，…，
min{d?n/2」-1，d?n/2」-1+?n/2」}，min{dn-1}}

(5)

產(chǎn)生兩個(gè)子向量l1和l2需要經(jīng)過「n/2?(符號「?表示向上取整)次比較器的使用，在CUDA中這「n/2?次比較可以同時(shí)執(zhí)行，對于要尋找到本輪迭代中的最小值來說，它必定存在于子向量l2中，循環(huán)迭代上一步的過程，繼續(xù)將向量l2一分為二輸入比較器，直至l2中只存在一個(gè)元素時(shí)結(jié)束迭代，此時(shí)最后的這個(gè)元素就是本輪迭代中得到的最小值.經(jīng)過上述的第一輪最小值選擇操作之后，距離向量變?yōu)閧e0，e1，…en-2，en-1}，并且en-1這個(gè)元素是該向量中的最小值，也是算法選擇出的第一個(gè)最近鄰元素，然后繼續(xù)對{e0，e1，…en-2}向量做如上操作，就能得到第二個(gè)最近鄰元素.經(jīng)過k次最小值查找之后，就可以查找出初始距離向量中的k個(gè)最近鄰元素，由于每輪迭代中可以在CUDA中開啟多線程并行查找，將原本時(shí)間復(fù)雜度為O(n)的比較階段降為O(1)，所以一次最小值查找的時(shí)間復(fù)雜度就等于迭代次數(shù)O(「log2n?)，那么整體k次最小值查找的時(shí)間復(fù)雜度就是O(k「log2n?)，相比于其它排序算法擁有更低的時(shí)間復(fù)雜度.

綜上所述，基于k次最小值查找的最近鄰選擇算法的詳細(xì)步驟如下：

1)將全局存儲器中的m*n的距離矩陣拷貝至共享內(nèi)存.

2)當(dāng)n的值小于1024時(shí)，在CUDA中開啟m個(gè)線程塊，每一個(gè)線程塊負(fù)責(zé)一個(gè)測試數(shù)據(jù)點(diǎn)距離向量的排序，每個(gè)線程塊內(nèi)部分配「n/2?個(gè)線程，每個(gè)線程負(fù)責(zé)兩個(gè)對應(yīng)位置上數(shù)據(jù)的比較器選擇，留下值較小的數(shù)據(jù)并更新到共享內(nèi)存上.

3)每一輪迭代通過syncthreads()函數(shù)同步塊內(nèi)的所有線程，同步完成之后繼續(xù)對子向量l2進(jìn)行比較器選擇，在完成O(「log2n?)次迭代之后，每個(gè)線程塊零號線程對應(yīng)位置上的共享內(nèi)存存儲的數(shù)據(jù)就是本輪最小值.

4)經(jīng)過k次選擇就得到屬于該測試數(shù)據(jù)的k個(gè)最近鄰.

當(dāng)n的值較大時(shí)，由于CUDA架構(gòu)的硬件限制，則每個(gè)向量需要使用「n/1024?個(gè)線程塊負(fù)責(zé)計(jì)算，通過找出這些線程塊輸出數(shù)據(jù)中的最小值從而得到最近鄰數(shù)據(jù).

3.3 基于雙調(diào)排序的最近鄰選擇

通知分析可知：基于k次最小值查找的最近鄰選擇算法，當(dāng)k值比較大時(shí)，經(jīng)過k次查找所需要的時(shí)間也越大.從時(shí)間復(fù)雜度上分析，如果k值接近于訓(xùn)練集數(shù)據(jù)量n時(shí)，該方法的時(shí)間復(fù)雜度就變?yōu)镺(nlogn)，與快速排序的時(shí)間復(fù)雜度相同，相當(dāng)于將距離向量完全排序，因此它不適應(yīng)k值較大的情況.本節(jié)介紹了一種基于CUDA的雙調(diào)排序算法，它同樣能夠利用多線程特性來并行執(zhí)行，適合于k值比較大的排序場合.

本算法定義一個(gè)序列s=(a0，a1，…，an-1)，如果滿足條件a0≤a1≤…≤an/2-1并且an/2≥an/2+1≥…≥an-1，那么稱序列s就是一個(gè)雙調(diào)序列.算法將任意一個(gè)長度為n(n為偶數(shù))的雙調(diào)序列s分為相等長度的兩段s1和s2，將s1中的元素和s2中的元素一一比較，即a[i]和a[i+n/2](i

經(jīng)過距離計(jì)算得到的m*n的距離矩陣，使用基于雙調(diào)排序的最近鄰選擇算法的詳細(xì)步驟如下：

1)對于每一個(gè)測試數(shù)據(jù)對應(yīng)的距離向量l，其長度為n，算法中雙調(diào)排序輸入的序列長度為2a(a>0)，當(dāng)n的大小不足2a時(shí)，補(bǔ)足最少個(gè)數(shù)的元素使n=2a，補(bǔ)足的元素統(tǒng)一為能夠表示的最大數(shù).

2)將數(shù)據(jù)拷貝至共享內(nèi)存，同時(shí)開啟若干個(gè)線程塊負(fù)責(zé)距離序列的雙調(diào)排序，塊內(nèi)總共進(jìn)行a輪迭代，前a-1輪進(jìn)行相鄰兩個(gè)單調(diào)性相反的序列合并，并分別按相反單調(diào)性遞歸進(jìn)行雙調(diào)排序.

3)第a輪時(shí)合并前面兩個(gè)長度為n/2的單調(diào)序列，開啟n/2個(gè)線程做兩兩元素的比較，按此方式迭代a輪，每輪都會開啟n/2個(gè)線程參與計(jì)算，直至最后一輪迭代時(shí)是長度為2的序列相比較，即可得出最后單調(diào)遞增的距離序列.

圖3 構(gòu)建雙調(diào)序列示意圖Fig.3 Diagram of a bitonic sequence

基于雙調(diào)排序的最近鄰選擇算法的時(shí)間復(fù)雜度為O((logn)2)，所以當(dāng)k較大時(shí)，該方法的時(shí)間復(fù)雜度是優(yōu)于基于k次最小值查找的最近鄰選擇算法.

3.4 基于CUDA的決定分類標(biāo)號

基于CUDA的決定分類標(biāo)號的具體步驟如下：在CUDA中開啟m個(gè)線程塊，每個(gè)線程塊負(fù)責(zé)一個(gè)測試數(shù)據(jù)分類標(biāo)號的統(tǒng)計(jì)，在共享內(nèi)存上開辟一個(gè)數(shù)組，數(shù)組長度為訓(xùn)練數(shù)據(jù)集的類別個(gè)數(shù).塊內(nèi)開啟k個(gè)線程，每個(gè)線程統(tǒng)計(jì)自己對應(yīng)位置上數(shù)據(jù)的類別標(biāo)號，采用CUDA內(nèi)部的原子加法操作統(tǒng)計(jì)，即atomicAdd()操作，在類別數(shù)組上取數(shù)據(jù)并執(zhí)行加法操作.最后同步線程塊內(nèi)線程，得到最終記錄數(shù)組，數(shù)組當(dāng)中數(shù)字最大的類別即為該測試集最終的分類.

3.5 GS_KNN算法的詳細(xì)步驟

綜合前面各節(jié)所述，GS_KNN算法的詳細(xì)步驟如算法2所示，算法流程如圖4所示.

圖4 GS_KNN算法流程圖Fig.4 Flow chart of GS_KNN algorithm

算法2.基于CUDA的GS_KNN算法

輸入：m個(gè)點(diǎn)的測試數(shù)據(jù)集，n個(gè)點(diǎn)的訓(xùn)練數(shù)據(jù)集，以及最近鄰個(gè)數(shù)k.

輸出：m個(gè)測試數(shù)據(jù)的分類結(jié)果

1)在CPU主機(jī)端初始化測試數(shù)據(jù)集，訓(xùn)練數(shù)據(jù)集和最近鄰個(gè)數(shù)k的內(nèi)存空間.

2)在GPU設(shè)備端分配測試數(shù)據(jù)集，訓(xùn)練數(shù)據(jù)集以及最近鄰個(gè)數(shù)k的內(nèi)存空間.

3)將主機(jī)端數(shù)據(jù)拷貝至設(shè)備端.

4)利用CUDA調(diào)用基于Cublas通用矩陣乘的距離計(jì)算內(nèi)核函數(shù)，得到所有測試數(shù)據(jù)集數(shù)據(jù)到訓(xùn)練數(shù)據(jù)集的距離矩陣.

5)隨機(jī)抽取一個(gè)測試數(shù)據(jù)的距離向量，分別進(jìn)行基于k次最小值查找和基于雙調(diào)排序的最近鄰選擇排序算法的實(shí)驗(yàn)，得到時(shí)間T1和T2.

6)如果T1小于T2，則選擇基于k次最小值查找的最近鄰選擇算法，進(jìn)行距離排序.否則，選擇基于雙調(diào)排序的最近鄰選擇算法來排序.

7)進(jìn)行基于CUDA的決定分類標(biāo)號，得到每個(gè)測試數(shù)據(jù)的分類結(jié)果.

8)將最后的分類結(jié)果從GPU設(shè)備端返回拷貝至CPU主機(jī)端并輸出.

4 實(shí)驗(yàn)與分析

4.1 實(shí)驗(yàn)環(huán)境

本文所使用的實(shí)驗(yàn)平臺為Intel Core i5-4590 CPU，主頻3.30GHz，內(nèi)存8GB，GPU為NVIDIA GTX750 Ti，計(jì)算能力為5.0，擁有5個(gè)流多處理器，640個(gè)SP.軟件環(huán)境為Windows 10，Visual studio 2010和CUDA7.0.

4.2 實(shí)驗(yàn)結(jié)果

本文實(shí)驗(yàn)所使用的第一個(gè)數(shù)據(jù)集為程序生成的模擬數(shù)據(jù)，可以指定測試集大小、訓(xùn)練集大小、每個(gè)樣本的類別以及它的屬性維度.第二個(gè)數(shù)據(jù)集為KDDCUP1999提供的網(wǎng)絡(luò)入侵檢測數(shù)據(jù)集，該數(shù)據(jù)集包含大量真實(shí)的網(wǎng)絡(luò)入侵?jǐn)?shù)據(jù)，數(shù)據(jù)集中每個(gè)元素包含41個(gè)維度的屬性，每個(gè)維度屬性為浮點(diǎn)類型的數(shù)據(jù)，總共有24種攻擊類型.由于每個(gè)維度上的度量單位不同，需要對輸入數(shù)據(jù)進(jìn)行標(biāo)準(zhǔn)化和歸一化，本文使用Wang等[15]的方法進(jìn)行數(shù)據(jù)的預(yù)處理.

4.2.1 固定m，k，改變n，d對加速的影響

本實(shí)驗(yàn)采用模擬數(shù)據(jù)，設(shè)定測試集的大小m為1200，訓(xùn)練集n的大小為210到215，這樣可以方便線程的訪問，測試集與訓(xùn)練集的維度d相同，k的值設(shè)為25.由于k值較小，所以經(jīng)過評估后算法采用基于k次最小值查找的最近鄰選擇.選擇了兩種算法來與GS_KNN算法作對比.第一種算法選用了基于CPU的超快近似最近鄰搜索算法庫，即ANN-C++[16]，另一種選用了BF-CUDA算法.改變n和d的大小，實(shí)驗(yàn)結(jié)果如表1所示，可以看出GS_KNN算法在相同數(shù)據(jù)集上的性能上要優(yōu)于ANN-C++算法和BF-CUDA算法，并且隨著訓(xùn)練集大小和維度的增加，GS_KNN算法相較于這兩個(gè)算法的加速比也越來越高，當(dāng)n等于215，d等于256時(shí)，GS_KNN算法相對于ANN-C++算法的加速比為145，相對于BF-CUDA算法的加速比為2.8.

4.2.2 固定m，n，d，改變k對排序的影響

實(shí)驗(yàn)(1)是在固定k=25的情況下進(jìn)行的，為了比較k值的大小對算法距離排序階段的影響，選取測試數(shù)據(jù)集大小為1200，測試數(shù)據(jù)集為16384，數(shù)據(jù)維度為32，第一次實(shí)驗(yàn)中GS_KNN算法的距離排序采用基于k次最小值查找的最近鄰選擇，第二次采用基于雙調(diào)排序的最近鄰選擇，然后每次實(shí)驗(yàn)都不斷增加k的值，觀察記錄距離排序階段的執(zhí)行時(shí)間，實(shí)驗(yàn)結(jié)果如圖5所示，可以看出當(dāng)k的值為46時(shí)，基于雙調(diào)排序和基于k次最小值查找的距離排序所用時(shí)間相同，當(dāng)k的值小于46時(shí)，k次最小值查找所需的時(shí)間要少于雙調(diào)排序.隨著k的不斷增加，雙調(diào)排序所花費(fèi)的時(shí)間不變，因?yàn)椴还躪值如何變化，算法都會得到一個(gè)完整的排序結(jié)果，而k次最小值查找的時(shí)間開銷呈線性增長趨勢，因?yàn)樗惴ǖ牡螖?shù)會隨著k的增加而線性增加，此時(shí)雙調(diào)排序的優(yōu)勢會越來越明顯.

表1 改變n和d時(shí)各算法運(yùn)行時(shí)間比較(單位：s)Table 1 Comparison of running time of each algorithm when changing n and d (unit：s)

圖5 改變k對距離排序時(shí)間的影響Fig.5 Effect of changing k on distance sorting time

4.2.3 算法準(zhǔn)確性驗(yàn)證

從KDDCUP99數(shù)據(jù)集中選取5000條數(shù)據(jù)，其中4900條正常數(shù)據(jù)，100條包含攻擊的入侵?jǐn)?shù)據(jù)作為訓(xùn)練集，再選500條包含正常數(shù)據(jù)以及訓(xùn)練集中未知的入侵?jǐn)?shù)據(jù)集作為測試集，分別運(yùn)用經(jīng)典的基于CPU的K最近鄰入侵檢測算法和GS_KNN算法，進(jìn)行10次實(shí)驗(yàn)取檢測率和誤報(bào)率的平均值，每次調(diào)整k值的大小，統(tǒng)計(jì)結(jié)果如表2所示.

表2 入侵檢測魯棒性檢測結(jié)果Table 2 Intrusion detection robustness test results

從結(jié)果來看，GS_KNN算法在檢測4類攻擊上與標(biāo)準(zhǔn)CPU版本的K最近鄰算法相比，檢測率和誤報(bào)率幾乎相同.造成略微差異的主要是由于在距離排序階段，離測試數(shù)據(jù)距離完全相等而類別標(biāo)號不同的訓(xùn)練數(shù)據(jù)被k的不同切分所導(dǎo)致.

5 結(jié) 論

本文在經(jīng)典串行KNN算法的基礎(chǔ)上，利用GPU的并行計(jì)算能力，提出了一種基于CUDA的并行KNN算法，即GS_KNN算法，該算法能有效處理大規(guī)模數(shù)據(jù)情況下的數(shù)據(jù)分類問題.算法利用矩陣乘的思想加速了距離矩陣的運(yùn)算，提出基于k次最小值查找和雙調(diào)排序兩種策略優(yōu)化距離排序，并在CUDA中完成分類標(biāo)號的統(tǒng)計(jì).它與經(jīng)典的BF-CUDA算法相比獲得2.8倍的加速比.然而本文在GPU和CPU之間的數(shù)據(jù)傳輸以及多GPU協(xié)同合作還有待研究，因此下一步盡可能優(yōu)化數(shù)據(jù)的存儲，利用GPU集群提升算法的性能.