• <tr id="yyy80"></tr>
  • <sup id="yyy80"></sup>
  • <tfoot id="yyy80"><noscript id="yyy80"></noscript></tfoot>
  • 99热精品在线国产_美女午夜性视频免费_国产精品国产高清国产av_av欧美777_自拍偷自拍亚洲精品老妇_亚洲熟女精品中文字幕_www日本黄色视频网_国产精品野战在线观看 ?

    基于異構(gòu)平臺的圖像中值濾波的OpenCL加速算法

    2024-03-15 06:47:40肖詩洋王鐳杜瑩肖漢
    關(guān)鍵詞:并行算法中值工作組

    肖詩洋,王鐳,杜瑩,肖漢

    (1.東南大學(xué) 土木工程學(xué)院,江蘇 南京 211189;2.鄭州師范學(xué)院 信息科學(xué)與技術(shù)學(xué)院,河南 鄭州 450044;3.鄭州師范學(xué)院 地理與旅游學(xué)院,河南 鄭州 450044)

    外界環(huán)境、傳感器件的質(zhì)量等各種因素會對圖像在采集、傳輸和存儲的過程中產(chǎn)生干擾.引入的噪聲致使圖像質(zhì)量下降,會嚴重影響圖像后續(xù)處理的效果.因此,為抑制圖像的噪聲,在圖像預(yù)處理中需要對圖像去噪.噪聲消除是蜂窩移動通信、圖像處理、雷達、聲納及其他應(yīng)用中需要解決的主要問題之一[1].

    目前圖像去噪方法有很多種,圖像的噪聲可根據(jù)噪聲類型選用空間域方法或變換域方法進行去噪,如中值濾波器、形態(tài)濾波器、高斯濾波器等.基于排序統(tǒng)計理論的中值濾波器是一種非線性數(shù)字濾波器,常用于消除圖像中的脈沖噪聲或椒鹽噪聲.語音、圖像和數(shù)字信號處理中的平滑操作可用中值濾波實現(xiàn).非線性中值濾波器既可以有效地消除脈沖干擾和隨機噪聲,又可實現(xiàn)對圖像的邊緣信息的充分保留.在大數(shù)據(jù)技術(shù)的推動下,數(shù)字圖像分辨率快速提高,對圖像處理算法的實時性和處理效率都提出了較高的要求.因此,探索高效快速的圖像處理方法和理論,是提高應(yīng)用系統(tǒng)能力的重要需求,也是對圖像處理技術(shù)發(fā)展提出的科學(xué)難題之一[2].

    利用現(xiàn)場可編程門陣列(field programmable gate array,FPGA)或數(shù)字信號處理器(digital signal processor,DSP)實現(xiàn)實時圖像中值濾波處理,要求增添額外的硬件,系統(tǒng)的復(fù)雜度和費用會增加.隨著將多核CPU處理單元引入到并行系統(tǒng),需要對多核CPU中的每個CPU內(nèi)核進行編程.該方法操作繁瑣,編程量大.近年來,在科學(xué)計算和工程計算領(lǐng)域中的加速已廣泛運用圖形處理器(graphics processing unit, GPU),較于CPU,GPU提供了更強勁的計算能力和更大的數(shù)據(jù)讀寫帶寬.統(tǒng)一計算設(shè)備架構(gòu)(compute unified device architecture,CUDA)使研究人員可以對GPU進行更高效的編程,GPU的軟件生態(tài)環(huán)境得到了質(zhì)的提升.但CUDA系統(tǒng)只適用于NVIDIA系列的GPU.面向異構(gòu)系統(tǒng)結(jié)構(gòu)的開放、免費的并行編程標準開放式計算語言(open computing language,OpenCL)則可適用于任意一種并行處理器.OpenCL支持CPU、GPU、FPGA等并行處理器,擁有獨特的任務(wù)并行執(zhí)行模型,特別有利于支持異構(gòu)計算[3].

    圖像中值濾波算法是一種計算密集型算法,算法的中值濾波處理對算法的性能造成很大的影響.隨著圖像像素數(shù)據(jù)量的增大,算法的計算量迅速增長,進而大幅增加了算法處理時間,算法性能成為亟待解決的問題.而異構(gòu)計算成為改善算法性能的一種很好的解決方案.本研究將采用OpenCL編程模型,圍繞單GPU對中值濾波算法的并行化、算法的深度優(yōu)化和多種并行計算方案的比對研究等問題展開.

    1) 提出了一種基于OpenCL的圖像中值濾波并行算法.該算法既具有29.74倍加速比的高速去噪功能,又能夠很好地在不同GPU平臺中實現(xiàn)可移植性.

    2) 采用了多元化的性能對比標準.本文在多種并行計算平臺上實現(xiàn)了對圖像中值濾波的處理,測試了3種并行方式對算法性能的影響.將基于OpenCL的中值濾波(OpenCL-based median filtering, OCL_MF)并行算法、基于OpenMP的中值濾波(OpenMP-based median filtering, OMP_MF)并行算法、基于CUDA的中值濾波(CUDA-based median filtering, CUDA_MF)并行算法以及相關(guān)文獻的算法進行性能比較,從橫縱2個方向的性能比較看出,OCL_MF并行算法均取得了較好的性能提升.

    1 相關(guān)研究

    經(jīng)過多年的不懈努力,圖像中值濾波算法在算法優(yōu)化上取得了不少的成果.Zhao等[4]通過改進的加權(quán)中值濾波器,在GPU上實現(xiàn)了快速的紋理過濾.PANG等[5]提出了一種自適應(yīng)中值濾波器去除圖像中的隨機值脈沖噪聲,易于并行處理的實現(xiàn).Liu等[6]通過并行隨機電路的硬件實現(xiàn)了大量隨機數(shù)排列,提高了中值濾波處理效率.Cadenas等[7]研究了一種快速尋找中位數(shù)的方法,用于中值濾波算法,提高了降噪效率.Oded[8]設(shè)計了基于GPU直方圖運算的中值濾波軟件并行可擴展算法,提高速度60倍.李余錢等[9]通過基于FPGA的中值濾波算法確定自適應(yīng)閾值來實現(xiàn)Sobel邊緣檢測的方法,提高了系統(tǒng)的實時性.Sachin等[10]提出了一種基于single rogram multiple data (SPMD)的分時處理的中值濾波算法,相比基于remote method invocation (RMI)方式的分布系統(tǒng)性能更佳.劉佳等[11]提出了一種多處理器NoC結(jié)構(gòu)的中值濾波算法,處理速度提高了3.6倍.陳思潤等[12]使用ARM硬件架構(gòu)的Cortex-A處理器,運用SIMD數(shù)據(jù)級并行計算技術(shù)NEON實現(xiàn)了中值濾波并行算法,速度提高了17倍.Zhao等[13]提出了一種基于GPU的實時加權(quán)中值濾波器,提高了算法的執(zhí)行效率.Mao等[14]設(shè)計了基于GPU的加權(quán)中值濾波算法,應(yīng)用于油氣勘探數(shù)據(jù)處理平臺,獲得了4倍加速比.廖文獻等[15]在Cortex嵌入式多處理器系統(tǒng)上進行了圖像中值濾波算法的并行化,速度提高了近4倍.Mursaev等[16]設(shè)計了一種基于FPGA的二維中值濾波器,提高了處理速度.張怡卓等[17]設(shè)計了基于FPGA硬件的并行中值濾波算法,獲得了3.68倍加速比.Vivek等[18]提出了基于具有深度流水線的FPGA的中值濾波算法,提高了去噪速度.

    總之,上述研究工作在中值濾波算法方面取得了很多進展.有些通過改進中值濾波算法來提高去噪速度,有些提出了基于FPGA、ARM和GPU等計算平臺的并行算法,應(yīng)用于實際系統(tǒng)中提高處理速度,有些是在特定硬件平臺上加速中值濾波算法的處理效率.然而,各種加速算法的性能提高效果不理想,特別是這些加速算法都不具有平臺獨立性特點,只能在特定的硬件平臺或應(yīng)用中實現(xiàn)性能提升.本文擬分析圖像去噪中經(jīng)典的中值濾波算法,利用熱點分析法將耗時最長的功能模塊從中值濾波算法中找出.從而使用OpenCL對這些功能模塊進行GPU并行化.借助GPU在大規(guī)模并行運算上的優(yōu)勢,實現(xiàn)對中值濾波算法的CPU+GPU協(xié)同計算,有效縮短算法執(zhí)行時間和實現(xiàn)算法在不同GPU計算平臺上的性能移植.

    2 算法的研究與分析

    2.1 算法基本過程

    中值濾波器作為一種非線性濾波器,是一種基于排序統(tǒng)計理論的有效抑制噪聲的非線性信號處理技術(shù)[19].1970年Tukey提出了中值濾波器并在一維信號處理技術(shù)中應(yīng)用,后又被二維圖像信號處理技術(shù)采用.在一定條件下,它可較好地克服平均值濾波、最小均方濾波等線性濾波器所產(chǎn)生的圖像細節(jié)模糊問題,對降低脈沖干擾及圖像掃描噪聲非常有效[20].

    中值濾波算法是指用像素點鄰域內(nèi)的所有像素點的灰度值的中位數(shù)代替像素點的灰度值,為了保證取中位數(shù)的便利性,鄰域內(nèi)的像素點的數(shù)目必須是正奇數(shù),可用式(1)來表示.

    (1)

    圖1 3×3濾波窗口Fig.1 3×3 filtering window

    中值濾波的具體實現(xiàn)步驟如下:

    1)在圖像中遍歷所有像素點,并以每個像素點為中心形成一系列中值濾波窗口;

    2)將濾波窗口內(nèi)的各個像素的灰度值讀取出來;

    3)把這些灰度值進行大小排序;

    4)將這些灰度值的中位數(shù)賦給對應(yīng)濾波窗口中心位置的像素,作為最終的輸出結(jié)果.

    2.2 算法熱點分析

    對于像幅大小為height×width的圖像進行的中值濾波算法主要包括以下模塊:1) 擴充圖像功能.拓展過程要求在左右方向進行原圖像拓展和上下方向拓展原圖像.該過程的時間復(fù)雜度為O(height×width+height×MODELDIM/2+width×MODELDIM/2.2)圖像中值濾波功能.需要將圖像中所有像素點的鄰域進行中值濾波處理,需要執(zhí)行height×width次,其時間復(fù)雜度為O(height×width).同時提取濾波窗口覆蓋下圖像子塊中像素點的灰度值,其時間復(fù)雜度為O(MODELDIM2/4).在圖像子塊中的灰度值進行排序,時間復(fù)雜度為O(MODELDIM4).由此可得,圖像中值濾波算法的時間復(fù)雜度為O(height×width×MODELDIM4),模塊2)是該算法的核心環(huán)節(jié).

    為實現(xiàn)圖像中值濾波算法的并行化,首先針對基于CPU的中值濾波(CPU-based median filtering, CPU_MF)算法進行熱點分析,定位算法的耗時步驟.CPU_MF算法在進行熱點分析時,是在像幅大小為5 471×5 682的圖像上進行測試,采用bmp圖像格式.將clock()函數(shù)插入到CPU_MF算法各主要步驟之間進行時間測試,獲得各主要步驟運算前后的時間點值,通過計算得到CPU_MF算法主要步驟的時間統(tǒng)計和整個算法的運行時間.

    從表1中可以看出,圖像中值濾波算法耗時主要集中在第2步.該步驟運行時間占串行算法總執(zhí)行時間的百分比為90.47%,其他函數(shù)處理時間只占據(jù)了不到1 s的時間,說明圖像中值濾波步驟是串行算法的性能瓶頸所在.如果采取一定并行措施大幅降低中值濾波模塊的處理時間,算法就可以獲得良好的加速效果.

    表1 圖像中值濾波算法各步驟運行時間及占比

    2.3 算法可并行化

    算法自身存在的任務(wù)依賴性影響著算法的可并行性.關(guān)聯(lián)性越低的任務(wù),其并行效果越好,反之則效果越差.中值濾波計算過程需要改進,以適合GPU的并行計算架構(gòu)的特點,并能夠充分利用GPU的各種具有各自優(yōu)勢的存儲器.本文在圖像邊界處理時,將待處理的原圖像矩陣行和列各補MODELDIM/2行和MODELDIM/2列.擴展處理圖像后,所有待處理像素矩陣均處于擴展圖像像素矩陣的內(nèi)部.此時對于每個像素點的處理均可采用一致的算法實現(xiàn),避免了分支處理,實現(xiàn)過程的高度統(tǒng)一得以保證,算法的并行潛力進一步提高.因此,可以分析得出圖像中值濾波功能具有三級并行性:

    1)像素級并行:椒鹽噪聲是以黑白點的形式疊加在圖像上,去噪則要對圖像中每個像素點的灰度值進行變換處理,這種任務(wù)具有相互獨立適于并行的特點.

    2)圖像區(qū)域中值排序級并行:像素點鄰域中圖像數(shù)據(jù)排序和計算中值任務(wù)相互獨立,適于并行執(zhí)行.

    3)窗口級并行:對每個濾波窗口覆蓋下的圖像域處理任務(wù)相互獨立,適于并行處理.

    中值濾波算法的并行性如圖2所示.圖2a中綠色區(qū)域為原始圖像待處理區(qū)域,局部處理時鄰域大小為3×3,白色區(qū)域為擴充圖像區(qū)域.像素點用方塊代表.在中值圖像濾波處理中,若對像素點A進行處理,即將子圖像塊1中的數(shù)據(jù)提取出來,進行數(shù)據(jù)排序和A像素點的灰度值處理等操作.然后依次對像素點B、C…X、Y進行處理.由于一系列子圖像區(qū)域之間沒有相互依賴性,所以,中值圖像濾波算法適于并行化處理.

    a.擴充圖像;b.像素點鄰域

    由此可見,圖像中值濾波算法的計算量非常大并且又具備很強的并行性.因此,算法適合利用GPU計算平臺進行大規(guī)模并行計算.針對以上分析,本文設(shè)計實現(xiàn)了圖像中值濾波算法在OpenCL計算平臺上的naive版本.采用OpenCL的多工作項對圖像中值濾波功能部分進行并行計算,一個工作項負責計算一個像素點鄰域的中值,如果系統(tǒng)啟動w個工作項,則算法的計算時間復(fù)雜度降為O((height×width×MODELDIM4)/w).

    3 圖像中值濾波算法的OpenCL實現(xiàn)

    3.1 算法的加速策略

    3.1.1 GPU核函數(shù)的設(shè)計與實現(xiàn)

    設(shè)計內(nèi)核的索引空間,將工作組和工作項的大小確定下來,即設(shè)定處理單元和計算單元的大小進行并行計算.主要實現(xiàn)提取像素點鄰域灰度值,濾波窗口內(nèi)像素值的冒泡排序和計算鄰域中位數(shù)并更新至像素點.詳細設(shè)計如下:

    2) 在采用OpenCL并行設(shè)計時,采用2層并行模式.將整幅圖像在邏輯上劃分成個子圖像塊,網(wǎng)格中的一個工作組對應(yīng)處理圖像中的某個子圖像塊,即為粗粒度并行.NDRange索引空間中的工作項依據(jù)索引地址和濾波窗口大小,提取位于全局存儲器中相應(yīng)像素點鄰域的灰度值.鄰域像素值用冒泡法進行排序,計算得到鄰域像素點的中位數(shù),并將相應(yīng)更新的灰度值數(shù)據(jù)寫入主存空間的像素點,即為細粒度并行.圖像中值濾波并行設(shè)計如圖3所示.

    圖3 計算空間并行化劃分Fig.3 Parallel partition of computational space

    3) 粗粒度的工作組并行和細粒度的工作項并行組成了OCL_MF并行算法計算架構(gòu).粗粒度并行存在于同一N-NDRange索引空間內(nèi)眾多工作組之間,工作組之間無需進行數(shù)據(jù)交換和通信.這樣,系統(tǒng)獲得了可擴展性:由于在任意一個計算單元上都可以執(zhí)行工作組子任務(wù),所以,在核心數(shù)量不同的GPU上都能正常運行OCL_MF并行系統(tǒng).細粒度并行存在于同一工作組內(nèi)的工作項之間,同一工作組內(nèi)的工作項之間可以進行數(shù)據(jù)交換和通信.因此,工作組中的不同工作項可通過(work_group_barrier)進行同步處理.

    在OpenCL的執(zhí)行模型中,需要滿足以下3個方面的限制:1)一個計算單元能夠支持并行運行工作項總量的限制;2)一個計算單元能夠支持并行運行工作組總量的限制;3)一個工作組能夠支持并行運行工作項總量的限制.一個計算單元所能并行運行的工作項總數(shù)是確定的,為了獲得OCL_MF算法的最大并行計算效率,工作組中所包含的工作項數(shù)目的設(shè)計至關(guān)重要.在對OCL_MF并行算法進行GPU實現(xiàn)的過程中,采用計算能力為5.0的GPU.其每個計算單元最多可以激活的工作項數(shù)目為2 048.在warp中1次可調(diào)度32個工作項,因此,為了提高并行效率,1個工作組容納工作項數(shù)應(yīng)是32的倍數(shù).同時鑒于工作組中寄存器等資源的限制,設(shè)計工作組的大小為16×16=256.256小于工作組中對最大工作項數(shù)1 024的限制.每個計算單元將執(zhí)行2 048/256=8個工作組,小于在計算單元上至多激活工作組數(shù)32的限制.因此,GPU中的所有工作項和工作組都同時處于并行計算狀態(tài).

    公民性缺失抑或制度供給不足?—對我國社區(qū)參與困境的微觀解讀……………………………徐 林 徐 暢(32)

    3.1.2 工作項坐標的索引機制

    圖像中值濾波器是用待處理像素相鄰范圍內(nèi)的幾個像素的中位數(shù)來代替該點的像素值.對一個子圖像塊數(shù)據(jù)進行中值濾波處理時,對圖像塊中每個像素數(shù)據(jù)的處理是無任何數(shù)據(jù)相關(guān)的.因此,可以將對每個像素數(shù)據(jù)的處理映射到工作組的相應(yīng)工作項中.每個工作組中的工作項均可獲得二維ID索引位置,將相應(yīng)像素點的鄰域數(shù)據(jù)調(diào)度到處理單元中處理.像素點和工作項之間的位置映射關(guān)系,如下所示:

    workItemx=get_local_id(0)+get_group_id(0)*get_local_size(0),

    (2)

    workItemx=get_local_id(1)+get_group_id(1)*get_local_size(1),

    (3)

    其中,get_local_id(0)和get_local_local_id(1)分別為工作項在工作組中在x和y方向上的索引號.get_group_id(0)和get_group_id(1)分別為工作組在索引空間中在x和y方向上的索引號.get_local_size(0)和get_local_size(1)分別為工作組在x和y方向上的維度.

    3.2 優(yōu)化設(shè)計

    3.2.1 設(shè)備與主機間的通信優(yōu)化

    減少在設(shè)備和主機之間的數(shù)據(jù)傳輸是OCL_MF并行濾波算法的設(shè)計關(guān)鍵.由于設(shè)備與主機間的數(shù)據(jù)傳輸帶寬遠低于設(shè)備之間的數(shù)據(jù)傳輸帶寬,若設(shè)備和主機之間數(shù)據(jù)傳過于頻繁,系統(tǒng)將陷入傳輸瓶頸而不能充分發(fā)揮GPU的并行計算能力,計算效率降低.為了降低傳輸數(shù)據(jù)頻度,OCL_MF把中值濾波功能中的全部計算任務(wù)全部映射到GPU中.一方面減少了中間數(shù)據(jù)在設(shè)備和主機之間的傳入傳出操作,另一方面一次性在設(shè)備和主機之間進行圖像數(shù)據(jù)傳輸.這種設(shè)計大大縮減了傳輸次數(shù),GPU的計算能力得到充分利用,計算密集度提高.

    3.2.2 工作組資源配置優(yōu)化

    在OpenCL框架中,系統(tǒng)定期在不同warp間進行切換以進行工作項調(diào)度,以滿足對計算單元中計算資源的充分利用.同一指令是在同一warp中的32個工作項間一同執(zhí)行,因此,每個工作組中工作項數(shù)量應(yīng)是32的整數(shù)倍,每個維度上的大小將根據(jù)任務(wù)量的情況確定.表2顯示了像幅大小為3 241×3 753時,系統(tǒng)中設(shè)置不同數(shù)量工作項時的運行時間.由表2可見,采用16×16的工作組維度時,系統(tǒng)性能最佳.

    表2 工作組維度對運算速度的影響

    4 實驗環(huán)境與實驗設(shè)計

    4.1 實驗環(huán)境

    1)硬件平臺.選取2個不同計算平臺的目的是驗證性能可移植性.CPU系統(tǒng)均為具有六核心的Intel Core i7 8700,主頻為3.2 GHz,系統(tǒng)存儲器為16.0 GB.2種GPU計算平臺的性能參數(shù)如表3所示.

    表3 GPU計算平臺性能參數(shù)

    4.2 實驗結(jié)果

    開展多組數(shù)據(jù)的對比實驗,需要對圖像數(shù)據(jù)進行預(yù)處理,通過剪裁取得了圖像大小分別為640×480、561×762、1 354×1 785、2 265×2 746、3 241×3 753、5 471×5 682、7 215×7 634和8 146×8 357共8組實驗數(shù)據(jù).

    圖4a是原始圖像,圖4b-e分別為用CPU_MF、OMP_MF、CUDA_MF和OCL_MF系統(tǒng)進行處理后的圖像.

    a.原始噪聲圖像;b.CPU_MF效果圖;c.OMP_MF效果圖;d.CUDA_MF效果圖;e.OCL_MF效果圖

    為檢驗圖像中值濾波算法的性能本文共設(shè)計4組實驗:第1組運行CPU_MF串行中值濾波算法,第2組運行OMP_MF并行中值濾波算法,第3組運行CUDA_MF的并行中值濾波算法,第4組運行OCL_MF并行中值濾波算法.針對測試圖像,4組中值濾波系統(tǒng)被多次運行,計算出各組中值濾波算法的平均耗時,數(shù)值結(jié)果保留小數(shù)點后2位,串/并行中值濾波算法執(zhí)行時間對比如表4所示.

    表4 不同計算平臺下圖像中值濾波算法執(zhí)行時間

    定義加速比

    (4)

    其中,s為加速比,Ts是CPU_MF串行算法執(zhí)行耗時,Tp是并行算法執(zhí)行耗時.

    定義相對加速比1 OMP_MF算法運算耗時與基于NVIDIA GPU的OCL_MF算法運算耗時的比值.

    定義相對加速比2 CUDA_MF算法運算耗時與基于NVIDIA GPU的OCL_MF算法運算耗時的比值.

    為了客觀評估系統(tǒng)的性能,采用了加速比以反映并行算法相比CPU串行算法整體速度改善情況.為了測評不同并行算法的性能,采用了相對加速比1以反映基于NVIDIA GPU的OCL_MF算法相比OMP_MF算法的性能改善情況,相對加速比2以反映基于NVIDIA GPU的OCL_MF算法相比CUDA_MF算法的性能改善情況.具體如表5所示.

    表5 不同計算平臺下圖像中值濾波并行算法性能對比

    4.3 CPU與GPU結(jié)果一致性實驗

    縮短圖像處理時間是中值濾波并行處理的目的,以求得更高圖像降噪速度.但是,若以損失圖像質(zhì)量為代價,并行化處理就沒有了意義.下面進行圖像中值濾波效果分析.

    4.3.1 宏觀層面結(jié)果一致性

    如圖4可見,原始圖像是一組被椒鹽噪聲污染的圖像,圖像中均疊加有許多黑白噪聲點.原始圖像經(jīng)過串行/并行圖像中值濾波處理后,去掉了大部分的黑白噪聲點,各幅圖像均變得更加平滑,對于去除脈沖型加性噪聲效果很好.同時,圖像中值濾波算法采用CPU串行實現(xiàn),OpenMP多線程并行實現(xiàn),CUDA架構(gòu)下GPU實現(xiàn)和OpenCL異構(gòu)平臺實現(xiàn),四者的濾波變換效果一樣.

    4.3.2 微觀層面結(jié)果一致性

    采用圖像灰度值直方圖來表示經(jīng)過串行系統(tǒng)和并行系統(tǒng)處理后的圖像數(shù)據(jù),如圖5所示.對比分析可見,從微觀層面上看,中值圖像濾波經(jīng)過串行和并行處理后的圖像直方圖均一樣.

    a.串行處理結(jié)果;b.OpenMP處理結(jié)果;c.CUDA處理結(jié)果;d.OpenCL處理結(jié)果

    因此,從宏觀和微觀2個層面來看,中值圖像濾波串行算法和并行算法雖然在設(shè)計方法和執(zhí)行時間上不同,然而在圖像處理的結(jié)果仍保持一致,算法的正確性和可行性得到驗證.

    4.4 實驗結(jié)果分析

    4.4.1 不同計算平臺上中值濾波算法運算時間分析

    通過對圖6的分析可見,當計算規(guī)模較小時,中值濾波算法采用GPU并行計算加速效果較為明顯,獲得了29.74倍的最大加速比.如像幅大小為561×762時,CPU_MF算法計算時間為149.00 ms,OMP_MF算法計算時間為86.61 ms,CUDA_MF并行算法計算時間為5.21 ms,OCL_MF并行算法計算時間為5.01 ms,傳統(tǒng)串行方式和多核并行方式的耗時遠高于OpenCL并行加速計算耗時.當像幅較大時,CPU_MF算法計算時間呈現(xiàn)近乎直線的上升趨勢,耗時急劇增加.OMP_MF算法的計算時間表現(xiàn)出緩慢上升,而GPU并行算法的計算時間則表現(xiàn)出更為平緩的上升.同時,像幅大小超過5471×5682時,GPU并行算法耗時出現(xiàn)了較快的增長趨勢.

    圖6 中值濾波算法運算時間對比Fig.6 Comparison of operation time of median filtering algorithm

    GPU設(shè)備和CPU主機相互協(xié)作完成算法的處理過程,期間大量的數(shù)據(jù)需要交換.這種數(shù)據(jù)交換均由PCI-E總線負責完成.但是GPU內(nèi)部帶寬要遠高于PCI-E總線帶寬.因此,當通過PCI-E總線進行大量的圖像數(shù)據(jù)交換時,受限其速度,算法執(zhí)行時間變長,GPU加速的效果出現(xiàn)了減緩的趨勢.

    下面將文獻[17]與OCL_MF并行算法的加速效果進行比對.由于多數(shù)文獻是利用中值濾波算法進行各種應(yīng)用研究,很少專門針對中值濾波算法進行加速效果的研究.因此,無法直接進行加速效果的對比.根據(jù)文獻[17]中提供的測試數(shù)據(jù),當圖像大小為640×480時,文獻[17]中基于FPGA加速的中值濾波并行算法的運算時間為6.144 ms,獲得了3.68倍性能提高.而由表4和表5的測試結(jié)果可知,本文基于OpenCL加速的中值濾波并行算法的運算時間為3.77 ms,獲得了5.04倍加速比.因此,本文并行算法的運算時間比文獻[17]中的算法短,并取得了更好的加速性能.

    4.4.2 并行計算平臺上中值濾波算法加速比分析

    從圖7可以看出,當圖像像幅大小在3 241×3 753以內(nèi)時,OpenMP并行方式的加速比曲線斜率變化不大,而GPU并行方式的加速比曲線斜率也較小;當像幅大小從3 241×3 753擴展到5 471×5 682時,OpenMP加速方式的曲線斜率仍沒有明顯變化,而GPU加速方式的曲線則出現(xiàn)一個較陡峭的下降;當像幅大小超過5 471×5 682時,OpenMP加速方式的曲線斜率依然是平穩(wěn)的上升態(tài)勢,而GPU加速方式的曲線則呈現(xiàn)了緩慢下降的趨勢.因此,從圖7中可見,隨著像幅規(guī)模的增加,在像幅大小的各個區(qū)間內(nèi)GPU曲線斜率的變化都較OpenMP方式曲線斜率有明顯變化.

    圖7 中值濾波并行算法加速比趨勢Fig.7 Speedup trend diagram of median filtering parallel algorithm

    曲線斜率的大小,在一定程度上反映出數(shù)據(jù)規(guī)模與運算時間的關(guān)系,即數(shù)據(jù)規(guī)模相同時,曲線斜率越大,說明該計算方式的耗時變化越劇烈.當曲線負斜率較大時,數(shù)據(jù)規(guī)模稍微增加,就導(dǎo)致運算時間的急劇增加.這時數(shù)據(jù)規(guī)模與時間消耗的性價比較低,形成計算效率的低峰期,并且擴展性也較差.

    據(jù)上可知,基于GPU的中值濾波算法的擴展性不如OMP_MF并行算法,GPU加速方式表現(xiàn)的更容易形成計算瓶頸.然而,GPU擁有的更加豐富的并行計算資源帶來了巨大的加速優(yōu)勢.在圖像規(guī)模增大時,仍然具有遠高于OpenMP多核并行計算產(chǎn)生的加速效果.所以,GPU的中值濾波算法的性能更有優(yōu)勢.

    由圖7可見,OCL_MF并行算法的運算速度遠高于OMP_MF并行算法,且隨著圖像規(guī)模的不斷增大,速度差距有縮小的趨勢.這是由于CPU核心數(shù)較少所致,當CPU處于滿載情況下,性能提升空間有限.同時線程創(chuàng)建和調(diào)度也存在時間開銷.而在一定的計算量范圍內(nèi),OpenCL中每個工作項有大致相同的計算時間,運算時間的增加僅由于更多的工作項和工作組與硬件之間交互造成的必要時間消耗.

    圖8中相對加速比2表明,像幅較小時,OCL_MF并行算法與CUDA_MF并行算法性能接近.隨著圖像規(guī)模的增大,OCL_MF并行算法性能相比CUDA_MF并行算法稍快,最高獲得了1.15倍性能提升.因此,當像幅較大時,OpenCL并行算法有更大性能優(yōu)勢.相對加速比1顯示出OCL_MF并行算法與OMP_MF并行算法性能相比有較大提高,最高獲得了17.29倍性能提升.同時,隨著像幅的增大,兩類并行算法的性能差距有逐漸縮小的趨勢.當像幅較小時,OpenCL并行算法有更大性能優(yōu)勢.

    圖8 相對加速比趨勢Fig.8 Relative acceleration ratio trend graph

    4.4.3 基于OpenCL的中值濾波并行算法可移植性分析

    Radeon RX 470采用的AMD 2012年提出的GCN架構(gòu),GTX 1070為NVIDIA 2016年提出的Pascal架構(gòu).GTX 1070的基礎(chǔ)頻率是Radeon RX 470的1.63倍,而且前者使用的存儲器帶寬也比后者更寬,寄存器數(shù)量也更多.因此,本文采用的AMD GPU卡性能不如NVIDIA GPU卡,并行算法的性能受到了一定的影響,如圖7所示.然而,OpenCL加速的圖像中值濾波并行算法在2種GPU平臺上均獲得了近30倍加速效果.因此,并行算法在異質(zhì)GPU計算平臺上獲得了較好的性能可移植性,符合軟硬件實際情況.

    5 結(jié)束語

    利用數(shù)字圖像數(shù)據(jù)呈規(guī)則格網(wǎng)分布和易于并行處理的特點,針對圖像中值濾波處理算法核心部分進行了并行方案的設(shè)計與實現(xiàn).算法索引空間的維度能夠自適應(yīng)于圖像的規(guī)模,對工作組的資源配置和不同設(shè)備間的數(shù)據(jù)傳輸進行了優(yōu)化.通過為并行計算任務(wù)合理地分配計算單元和處理單元,充分挖掘了GPU的并行計算能力,圖像處理的效率得以提高.實驗結(jié)果顯示,在NVIDIA GTX 1070平臺上實現(xiàn)的OCL_MF并行算法與CPU_MF串行算法,OMP_MF和CUDA_MF兩種并行算法性能相比,加速比分別獲得了29.74倍、17.29倍和1.15倍,算法性能得到極大提升.同時,該OpenCL加速的并行算法在AMD Radeon RX 470和NVIDIA GTX 1070平臺上均獲得了相近的加速比,實現(xiàn)了在異質(zhì)GPU計算平臺間的性能移植.本文提出的基于OpenCL的圖像中值濾波并行處理方法能夠有效縮短系統(tǒng)運算時間,實時地完成較大像幅的圖像中值濾波處理,對其他圖像處理應(yīng)用也具有一定的借鑒意義.

    本文研究還有進一步優(yōu)化的空間,有待做更加深入的探索:擬在GPU集群上將MPI和OpenCL技術(shù)相結(jié)合,更大圖像塊之間的并行由MPI完成,圖像塊內(nèi)的并行由每個節(jié)點上的GPU完成.通過GPU集群將使系統(tǒng)的處理速度更快,以爭取在更加短的時間內(nèi)完成對更大尺寸圖像的處理工作.

    猜你喜歡
    并行算法中值工作組
    地圖線要素綜合化的簡遞歸并行算法
    肖幼率工作組赴戴家湖涵指導(dǎo)搶險
    治淮(2020年8期)2020-09-22 06:25:46
    Lagrange中值定理的巧妙應(yīng)用
    32個工作組印跡 >
    中國民政(2017年13期)2017-08-01 00:07:27
    微分中值定理教法研討
    基于GPU的GaBP并行算法研究
    后中值波電流脈沖MIG焊工藝
    磁縣政協(xié)專題聽取委員工作組2015年工作匯報
    鄉(xiāng)音(2016年2期)2016-02-26 20:38:40
    百項能效標準推進工程聯(lián)合工作組會議在京召開
    基于GPU的分類并行算法的研究與實現(xiàn)
    亚洲精品久久久久久婷婷小说| 国产不卡av网站在线观看| 韩国高清视频一区二区三区| 久久精品熟女亚洲av麻豆精品| 波野结衣二区三区在线| 精品国产乱码久久久久久小说| 成人漫画全彩无遮挡| 一本大道久久a久久精品| 亚洲精品国产av蜜桃| 亚洲,欧美精品.| 中文天堂在线官网| 国产精品人妻久久久影院| 99久久综合免费| 超碰97精品在线观看| 另类亚洲欧美激情| 只有这里有精品99| 我的亚洲天堂| 久久久久久久大尺度免费视频| 黄片小视频在线播放| 亚洲国产欧美一区二区综合| www.熟女人妻精品国产| 国产精品国产三级国产专区5o| 两个人免费观看高清视频| 久久精品国产亚洲av高清一级| 亚洲av男天堂| 色播在线永久视频| 国产一区有黄有色的免费视频| 国产高清不卡午夜福利| 十分钟在线观看高清视频www| 卡戴珊不雅视频在线播放| 欧美成人精品欧美一级黄| 人妻一区二区av| 亚洲精品视频女| 国产精品免费大片| 国产日韩欧美视频二区| 亚洲久久久国产精品| 亚洲国产欧美一区二区综合| 国产日韩一区二区三区精品不卡| 蜜桃国产av成人99| av有码第一页| 精品一区二区免费观看| 亚洲欧美清纯卡通| 精品一区二区三区四区五区乱码 | 亚洲伊人色综图| 999精品在线视频| 精品一品国产午夜福利视频| 欧美亚洲 丝袜 人妻 在线| 欧美日韩亚洲综合一区二区三区_| 中国国产av一级| 国产成人啪精品午夜网站| 99九九在线精品视频| www.自偷自拍.com| 黑丝袜美女国产一区| 国产精品秋霞免费鲁丝片| 卡戴珊不雅视频在线播放| 亚洲专区中文字幕在线 | 国产日韩一区二区三区精品不卡| 国产日韩欧美视频二区| 亚洲成av片中文字幕在线观看| 午夜av观看不卡| 综合色丁香网| 亚洲av成人不卡在线观看播放网 | 国产伦理片在线播放av一区| 观看av在线不卡| 久久精品久久精品一区二区三区| 亚洲精品一区蜜桃| 老司机影院毛片| 色婷婷av一区二区三区视频| 女人爽到高潮嗷嗷叫在线视频| 精品亚洲乱码少妇综合久久| 国产福利在线免费观看视频| 国产精品女同一区二区软件| 亚洲精品日本国产第一区| 久久久久久免费高清国产稀缺| 在线免费观看不下载黄p国产| 曰老女人黄片| 久久久久国产精品人妻一区二区| 亚洲三区欧美一区| 高清在线视频一区二区三区| 精品一区二区三卡| 精品人妻熟女毛片av久久网站| 亚洲激情五月婷婷啪啪| 校园人妻丝袜中文字幕| 精品亚洲乱码少妇综合久久| 午夜福利网站1000一区二区三区| 久久午夜综合久久蜜桃| 亚洲图色成人| 天天操日日干夜夜撸| 高清av免费在线| 一级a爱视频在线免费观看| 国产午夜精品一二区理论片| 观看av在线不卡| 91精品国产国语对白视频| 韩国精品一区二区三区| 亚洲成人av在线免费| 18在线观看网站| 精品一区二区三区四区五区乱码 | 97在线人人人人妻| 女人高潮潮喷娇喘18禁视频| 日韩 欧美 亚洲 中文字幕| 欧美老熟妇乱子伦牲交| 看十八女毛片水多多多| 七月丁香在线播放| 欧美 亚洲 国产 日韩一| 欧美97在线视频| 伊人亚洲综合成人网| 久久久久精品国产欧美久久久 | 9色porny在线观看| 久久精品国产a三级三级三级| 成人免费观看视频高清| 国产亚洲欧美精品永久| 精品少妇一区二区三区视频日本电影 | 久久av网站| 哪个播放器可以免费观看大片| 日日摸夜夜添夜夜爱| 新久久久久国产一级毛片| 男女之事视频高清在线观看 | 亚洲国产精品一区三区| 在线观看www视频免费| 亚洲精品美女久久av网站| 日韩不卡一区二区三区视频在线| 欧美国产精品va在线观看不卡| 成人午夜精彩视频在线观看| 成人国产麻豆网| 人妻人人澡人人爽人人| 亚洲人成网站在线观看播放| 免费在线观看黄色视频的| 王馨瑶露胸无遮挡在线观看| 午夜福利乱码中文字幕| 在线观看免费午夜福利视频| 老司机亚洲免费影院| 日韩欧美一区视频在线观看| 别揉我奶头~嗯~啊~动态视频 | 欧美久久黑人一区二区| 久久久久久久大尺度免费视频| 99国产综合亚洲精品| 啦啦啦在线观看免费高清www| 97精品久久久久久久久久精品| 人妻一区二区av| 一区二区av电影网| 国产精品99久久99久久久不卡 | netflix在线观看网站| 无遮挡黄片免费观看| 国产精品久久久久久人妻精品电影 | av在线老鸭窝| 日本av免费视频播放| 午夜福利视频精品| 国产精品香港三级国产av潘金莲 | 亚洲精品国产一区二区精华液| 精品卡一卡二卡四卡免费| 亚洲一区中文字幕在线| 秋霞在线观看毛片| 国产淫语在线视频| 国产av码专区亚洲av| 视频区图区小说| 久久久久久久久免费视频了| 最近的中文字幕免费完整| 国产精品香港三级国产av潘金莲 | 国产精品成人在线| av一本久久久久| 在线 av 中文字幕| 在线观看人妻少妇| 久久女婷五月综合色啪小说| 日韩伦理黄色片| 我的亚洲天堂| 在线观看www视频免费| 亚洲欧美成人精品一区二区| 18在线观看网站| 中文字幕色久视频| 亚洲国产中文字幕在线视频| xxxhd国产人妻xxx| 黄网站色视频无遮挡免费观看| 80岁老熟妇乱子伦牲交| 亚洲人成电影观看| 美女大奶头黄色视频| 国产一区二区在线观看av| 精品一区二区三区av网在线观看 | 国产1区2区3区精品| 亚洲欧美中文字幕日韩二区| 天天躁夜夜躁狠狠久久av| 老鸭窝网址在线观看| 色婷婷av一区二区三区视频| 国产一区二区 视频在线| 麻豆乱淫一区二区| 亚洲在久久综合| 91国产中文字幕| 777米奇影视久久| 欧美日韩国产mv在线观看视频| 热re99久久精品国产66热6| 人妻 亚洲 视频| 国产有黄有色有爽视频| av国产精品久久久久影院| 波多野结衣av一区二区av| 在线观看免费日韩欧美大片| 综合色丁香网| 久久精品久久久久久久性| 日日撸夜夜添| 亚洲成人av在线免费| 黄色视频不卡| av网站免费在线观看视频| 麻豆av在线久日| 成人国产麻豆网| 黑人欧美特级aaaaaa片| 国产精品免费视频内射| 亚洲国产看品久久| 国产伦人伦偷精品视频| 亚洲免费av在线视频| 欧美日韩av久久| 亚洲成人一二三区av| av视频免费观看在线观看| 中文字幕人妻丝袜制服| 99热国产这里只有精品6| 中文字幕高清在线视频| 另类亚洲欧美激情| 七月丁香在线播放| 18禁观看日本| 免费人妻精品一区二区三区视频| 国产在视频线精品| 91精品三级在线观看| 大码成人一级视频| 久久精品国产a三级三级三级| 亚洲精品自拍成人| 亚洲国产最新在线播放| 男女床上黄色一级片免费看| 高清av免费在线| 女性被躁到高潮视频| kizo精华| 国产麻豆69| 丝袜美足系列| 欧美黄色片欧美黄色片| 免费久久久久久久精品成人欧美视频| 亚洲成人免费av在线播放| 一区二区三区精品91| 日韩免费高清中文字幕av| 人妻一区二区av| 自拍欧美九色日韩亚洲蝌蚪91| 在线亚洲精品国产二区图片欧美| 亚洲人成电影观看| 曰老女人黄片| 亚洲成色77777| 国产片内射在线| 午夜福利视频在线观看免费| 看非洲黑人一级黄片| 国产精品免费视频内射| 男女下面插进去视频免费观看| 亚洲,一卡二卡三卡| 一区福利在线观看| 亚洲成国产人片在线观看| 欧美 日韩 精品 国产| 美女大奶头黄色视频| av国产久精品久网站免费入址| 久久天堂一区二区三区四区| 国产精品免费大片| 国产又爽黄色视频| 18禁裸乳无遮挡动漫免费视频| 欧美 日韩 精品 国产| 日韩av不卡免费在线播放| 国语对白做爰xxxⅹ性视频网站| 十八禁网站网址无遮挡| 丰满迷人的少妇在线观看| 天天躁狠狠躁夜夜躁狠狠躁| 久久久精品国产亚洲av高清涩受| 国产欧美日韩一区二区三区在线| 91aial.com中文字幕在线观看| av女优亚洲男人天堂| 免费高清在线观看日韩| 91老司机精品| 宅男免费午夜| 国产高清国产精品国产三级| 成年动漫av网址| 久久这里只有精品19| 国产无遮挡羞羞视频在线观看| 国产福利在线免费观看视频| 超碰成人久久| 精品国产露脸久久av麻豆| 国产av国产精品国产| 赤兔流量卡办理| xxx大片免费视频| 人人妻,人人澡人人爽秒播 | 亚洲精品视频女| av不卡在线播放| 国产免费现黄频在线看| av在线播放精品| 亚洲精品一二三| 又粗又硬又长又爽又黄的视频| 韩国精品一区二区三区| 国产一区二区三区av在线| 97精品久久久久久久久久精品| 久久人妻熟女aⅴ| 热re99久久国产66热| 亚洲精品美女久久久久99蜜臀 | 亚洲欧美一区二区三区国产| 在线看a的网站| 精品少妇内射三级| 国产一区二区激情短视频 | 丰满乱子伦码专区| 成年动漫av网址| 亚洲一级一片aⅴ在线观看| 自线自在国产av| 午夜av观看不卡| 一边摸一边做爽爽视频免费| 久久天躁狠狠躁夜夜2o2o | 美女中出高潮动态图| 桃花免费在线播放| 日日啪夜夜爽| 国产成人免费观看mmmm| 丁香六月天网| 最近最新中文字幕免费大全7| 午夜精品国产一区二区电影| 日本色播在线视频| 悠悠久久av| 日韩大码丰满熟妇| 日韩中文字幕欧美一区二区 | 精品一区在线观看国产| 欧美黑人精品巨大| 丰满乱子伦码专区| kizo精华| 久久热在线av| 亚洲久久久国产精品| 国产高清国产精品国产三级| 宅男免费午夜| 亚洲情色 制服丝袜| 考比视频在线观看| 国产高清国产精品国产三级| 欧美老熟妇乱子伦牲交| 久久综合国产亚洲精品| 国产成人一区二区在线| 丰满迷人的少妇在线观看| 国产又色又爽无遮挡免| 久久影院123| 午夜免费男女啪啪视频观看| 成人影院久久| 国产一卡二卡三卡精品 | 无遮挡黄片免费观看| 少妇人妻久久综合中文| 日本91视频免费播放| 亚洲国产av新网站| 亚洲伊人色综图| 午夜老司机福利片| 97在线人人人人妻| 国产成人欧美在线观看 | 国产精品久久久久久人妻精品电影 | www.av在线官网国产| 国产一区二区 视频在线| 国产不卡av网站在线观看| 欧美国产精品一级二级三级| kizo精华| 色视频在线一区二区三区| 一本一本久久a久久精品综合妖精| 成人影院久久| 中文字幕人妻丝袜一区二区 | 老司机影院毛片| 精品少妇内射三级| 精品国产一区二区三区久久久樱花| 午夜日韩欧美国产| 国产精品一区二区在线观看99| 久久精品久久精品一区二区三区| 女人久久www免费人成看片| 高清不卡的av网站| 久久国产精品大桥未久av| 9191精品国产免费久久| 免费观看av网站的网址| 日本欧美国产在线视频| 欧美日韩一区二区视频在线观看视频在线| 国语对白做爰xxxⅹ性视频网站| 国产av精品麻豆| 国产免费一区二区三区四区乱码| 国产av精品麻豆| 少妇猛男粗大的猛烈进出视频| 人妻 亚洲 视频| 少妇人妻久久综合中文| 中文字幕高清在线视频| 国产成人系列免费观看| 日韩av不卡免费在线播放| 亚洲国产欧美网| 秋霞伦理黄片| av免费观看日本| 久久久久网色| 国产探花极品一区二区| 最近最新中文字幕免费大全7| 青春草国产在线视频| 免费观看人在逋| 欧美日韩视频高清一区二区三区二| 十八禁高潮呻吟视频| 黄片无遮挡物在线观看| a级毛片黄视频| 女人久久www免费人成看片| 亚洲精品国产色婷婷电影| 涩涩av久久男人的天堂| 日本午夜av视频| 一级a爱视频在线免费观看| 精品国产乱码久久久久久男人| 久久精品久久精品一区二区三区| 另类亚洲欧美激情| 啦啦啦在线观看免费高清www| 黑人巨大精品欧美一区二区蜜桃| 亚洲av欧美aⅴ国产| 在线 av 中文字幕| 老鸭窝网址在线观看| 在线免费观看不下载黄p国产| 国产极品天堂在线| 另类精品久久| 久久精品久久久久久噜噜老黄| 国产在线一区二区三区精| 色吧在线观看| 久久精品熟女亚洲av麻豆精品| 黑人猛操日本美女一级片| 国产午夜精品一二区理论片| xxxhd国产人妻xxx| 色网站视频免费| 免费在线观看视频国产中文字幕亚洲 | 亚洲图色成人| 久久久久国产一级毛片高清牌| 在线观看免费视频网站a站| 国产片内射在线| 男女床上黄色一级片免费看| 热re99久久国产66热| 香蕉丝袜av| 中文乱码字字幕精品一区二区三区| 一级片免费观看大全| 在线免费观看不下载黄p国产| 青草久久国产| 老熟女久久久| 大码成人一级视频| 国产不卡av网站在线观看| 日韩成人av中文字幕在线观看| 五月开心婷婷网| 色播在线永久视频| 成人毛片60女人毛片免费| 亚洲精品第二区| 亚洲天堂av无毛| 黑人巨大精品欧美一区二区蜜桃| 大码成人一级视频| 宅男免费午夜| 91精品伊人久久大香线蕉| 香蕉丝袜av| 下体分泌物呈黄色| 亚洲精品,欧美精品| 中文字幕制服av| 少妇人妻久久综合中文| 国产色婷婷99| 人人澡人人妻人| 大香蕉久久网| 亚洲精品国产av成人精品| 在线看a的网站| 韩国av在线不卡| 中文字幕人妻熟女乱码| videos熟女内射| 亚洲综合色网址| 欧美老熟妇乱子伦牲交| 精品少妇内射三级| 国产成人精品久久二区二区91 | 精品免费久久久久久久清纯 | 亚洲久久久国产精品| 中文字幕高清在线视频| 国产成人精品久久二区二区91 | 欧美在线一区亚洲| 婷婷色av中文字幕| 色吧在线观看| 国产精品国产av在线观看| 亚洲精品久久午夜乱码| 欧美激情高清一区二区三区 | bbb黄色大片| 在线 av 中文字幕| 久久精品国产亚洲av涩爱| 中文字幕最新亚洲高清| 午夜福利乱码中文字幕| 18禁国产床啪视频网站| av福利片在线| 老汉色∧v一级毛片| 亚洲av中文av极速乱| 亚洲欧美中文字幕日韩二区| 男女免费视频国产| 国产精品久久久久成人av| 人体艺术视频欧美日本| 国产免费一区二区三区四区乱码| 日韩熟女老妇一区二区性免费视频| 日本爱情动作片www.在线观看| 亚洲,欧美,日韩| 亚洲国产精品一区三区| 欧美老熟妇乱子伦牲交| 在线观看国产h片| 国产不卡av网站在线观看| 精品一区在线观看国产| 另类亚洲欧美激情| 国产亚洲最大av| 80岁老熟妇乱子伦牲交| 精品一区在线观看国产| 国产成人系列免费观看| 日本vs欧美在线观看视频| 国产乱人偷精品视频| 别揉我奶头~嗯~啊~动态视频 | 免费观看a级毛片全部| 国产亚洲一区二区精品| 又大又爽又粗| 美女国产高潮福利片在线看| 国产成人精品久久二区二区91 | 80岁老熟妇乱子伦牲交| 国产精品欧美亚洲77777| 青春草视频在线免费观看| 啦啦啦在线观看免费高清www| 日本欧美视频一区| 可以免费在线观看a视频的电影网站 | 久久女婷五月综合色啪小说| 午夜日韩欧美国产| 成人亚洲欧美一区二区av| 精品久久久精品久久久| 国产av精品麻豆| 欧美 亚洲 国产 日韩一| 国产伦人伦偷精品视频| 天天躁日日躁夜夜躁夜夜| 女性生殖器流出的白浆| 亚洲,一卡二卡三卡| 伦理电影大哥的女人| 王馨瑶露胸无遮挡在线观看| 九草在线视频观看| 中文字幕av电影在线播放| 亚洲中文av在线| 99久久综合免费| 99国产综合亚洲精品| 欧美人与性动交α欧美软件| 欧美xxⅹ黑人| 欧美日韩av久久| 视频在线观看一区二区三区| 麻豆乱淫一区二区| 国产亚洲精品第一综合不卡| 这个男人来自地球电影免费观看 | 91精品三级在线观看| av在线播放精品| 久久精品人人爽人人爽视色| 欧美乱码精品一区二区三区| 天堂中文最新版在线下载| 老司机深夜福利视频在线观看 | 国产欧美日韩一区二区三区在线| 日韩精品有码人妻一区| 韩国精品一区二区三区| 中国三级夫妇交换| 精品国产一区二区三区四区第35| 国产精品亚洲av一区麻豆 | av国产精品久久久久影院| 日韩,欧美,国产一区二区三区| 日本av手机在线免费观看| 国产欧美亚洲国产| 王馨瑶露胸无遮挡在线观看| 国产有黄有色有爽视频| 在线 av 中文字幕| 亚洲av中文av极速乱| 久久精品久久久久久久性| 黄片无遮挡物在线观看| 国产精品蜜桃在线观看| 欧美日韩一级在线毛片| 99久久人妻综合| 18禁裸乳无遮挡动漫免费视频| 男人添女人高潮全过程视频| 男女高潮啪啪啪动态图| 高清视频免费观看一区二区| 久久青草综合色| 国产男女超爽视频在线观看| 超碰成人久久| 久久免费观看电影| 妹子高潮喷水视频| bbb黄色大片| 少妇精品久久久久久久| 国产精品亚洲av一区麻豆 | 国产精品三级大全| 国产又色又爽无遮挡免| 校园人妻丝袜中文字幕| 波野结衣二区三区在线| 在现免费观看毛片| 国产黄色免费在线视频| 亚洲国产欧美日韩在线播放| 精品一区二区免费观看| 日韩视频在线欧美| 制服人妻中文乱码| 在线天堂中文资源库| 91精品伊人久久大香线蕉| 中文天堂在线官网| 高清在线视频一区二区三区| 国产成人av激情在线播放| 人人妻人人澡人人爽人人夜夜| 2021少妇久久久久久久久久久| 国产探花极品一区二区| 亚洲色图 男人天堂 中文字幕| 国产一区亚洲一区在线观看| 亚洲图色成人| 久久久久国产一级毛片高清牌| 秋霞伦理黄片| 欧美日韩视频高清一区二区三区二| 在线天堂中文资源库| 日韩欧美精品免费久久| 人成视频在线观看免费观看| 亚洲精品在线美女| 亚洲国产av新网站| 美女中出高潮动态图| 国产精品女同一区二区软件| 久久久国产精品麻豆| 五月天丁香电影| 男女下面插进去视频免费观看| 欧美精品亚洲一区二区| 母亲3免费完整高清在线观看| 最近2019中文字幕mv第一页| 尾随美女入室| 又粗又硬又长又爽又黄的视频| 另类亚洲欧美激情| 一二三四中文在线观看免费高清| 狠狠婷婷综合久久久久久88av| 最近2019中文字幕mv第一页| 久久精品亚洲av国产电影网| 18禁观看日本| 可以免费在线观看a视频的电影网站 | 国产淫语在线视频| 久久人人爽av亚洲精品天堂| 女人久久www免费人成看片| 国产精品嫩草影院av在线观看| 亚洲精品自拍成人|