• <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噜噜一区二区三区| av卡一久久| 欧美日韩乱码在线| 男女做爰动态图高潮gif福利片| 国产三级在线视频| 三级经典国产精品| 国产激情偷乱视频一区二区| 中文在线观看免费www的网站| 97热精品久久久久久| 亚洲av熟女| 久久综合国产亚洲精品| 最好的美女福利视频网| 成人永久免费在线观看视频| 国产精品伦人一区二区| 国产视频内射| 麻豆国产av国片精品| 嫩草影院新地址| 国产视频首页在线观看| 精品久久久久久久久av| 国产黄a三级三级三级人| 深夜a级毛片| 丰满的人妻完整版| 日本一本二区三区精品| 国产在视频线在精品| 黄色欧美视频在线观看| 国产精品麻豆人妻色哟哟久久 | 亚洲国产精品合色在线| 欧美激情国产日韩精品一区| a级一级毛片免费在线观看| 国产v大片淫在线免费观看| 色播亚洲综合网| 淫秽高清视频在线观看| 精品久久久久久久久亚洲| 中国美白少妇内射xxxbb| 天堂网av新在线| av在线老鸭窝| 亚洲国产精品成人久久小说 | 亚洲久久久久久中文字幕| 国产精品爽爽va在线观看网站| 成人三级黄色视频| 国产成人精品久久久久久| 免费av不卡在线播放| 看免费成人av毛片| 两性午夜刺激爽爽歪歪视频在线观看| 亚洲欧美日韩东京热| 久久亚洲国产成人精品v| 18+在线观看网站| 成人午夜精彩视频在线观看| 国产精品不卡视频一区二区| 人妻制服诱惑在线中文字幕| 在线a可以看的网站| 亚洲在久久综合| 中文资源天堂在线| 美女被艹到高潮喷水动态| 欧美高清成人免费视频www| 免费一级毛片在线播放高清视频| 美女被艹到高潮喷水动态| 成人永久免费在线观看视频| .国产精品久久| 熟女电影av网| 亚洲真实伦在线观看| 亚洲欧美日韩无卡精品| 一级黄色大片毛片| 变态另类成人亚洲欧美熟女| 黄色一级大片看看| 九草在线视频观看| 九草在线视频观看| 白带黄色成豆腐渣| 亚洲国产精品合色在线| 亚洲欧美日韩东京热| 婷婷精品国产亚洲av| 成人永久免费在线观看视频| 久久鲁丝午夜福利片| kizo精华| 欧美一级a爱片免费观看看| 国产 一区 欧美 日韩| 国产在线男女| 观看免费一级毛片| 国产中年淑女户外野战色| 日韩亚洲欧美综合| 禁无遮挡网站| 亚洲最大成人手机在线| 日韩亚洲欧美综合| 99久久精品国产国产毛片| 99久久精品国产国产毛片| 高清毛片免费观看视频网站| а√天堂www在线а√下载| 久久久久久久午夜电影| 精品久久国产蜜桃| 国产一区二区激情短视频| 亚洲自偷自拍三级| 国产精品人妻久久久影院| 色吧在线观看| 99热网站在线观看| 久久精品国产自在天天线| 亚洲精品久久久久久婷婷小说 | 成人无遮挡网站| 国产亚洲av嫩草精品影院| 一本久久精品| 午夜福利视频1000在线观看| 99riav亚洲国产免费| 卡戴珊不雅视频在线播放| 亚洲av一区综合| 一级毛片久久久久久久久女| 国产视频首页在线观看| 亚洲国产欧美人成| 久久人人爽人人爽人人片va| 精品午夜福利在线看| 国产高清视频在线观看网站| 亚洲国产欧洲综合997久久,| 亚洲av.av天堂| 99久久精品热视频| avwww免费| 国产三级在线视频| 久久韩国三级中文字幕| av视频在线观看入口| 亚洲七黄色美女视频| 狠狠狠狠99中文字幕| 免费观看人在逋| 欧美zozozo另类| 欧美成人一区二区免费高清观看| 97热精品久久久久久| 国内精品宾馆在线| 99热全是精品| 亚洲精品456在线播放app| 欧美另类亚洲清纯唯美| 真实男女啪啪啪动态图| 午夜免费男女啪啪视频观看| 99热这里只有是精品50| 99久久精品一区二区三区| 免费在线观看成人毛片| 最近2019中文字幕mv第一页| 欧美成人a在线观看| 99riav亚洲国产免费| 久久九九热精品免费| 99在线视频只有这里精品首页| 美女 人体艺术 gogo| 中文字幕免费在线视频6| 又黄又爽又刺激的免费视频.| 啦啦啦观看免费观看视频高清| 日日摸夜夜添夜夜添av毛片| 亚洲国产精品合色在线| 99久久九九国产精品国产免费| 国产精品一区二区在线观看99 | 国产精品av视频在线免费观看| 成人国产麻豆网| 久久婷婷人人爽人人干人人爱| 久久这里只有精品中国| 国产精品人妻久久久影院| videossex国产| 狠狠狠狠99中文字幕| 狠狠狠狠99中文字幕| 久久久国产成人精品二区| 国内久久婷婷六月综合欲色啪| a级一级毛片免费在线观看| 成人欧美大片| 欧美人与善性xxx| 国产白丝娇喘喷水9色精品| 日本成人三级电影网站| 国内精品美女久久久久久| 欧美精品一区二区大全| 久久精品人妻少妇| 国产精品福利在线免费观看| 成人三级黄色视频| 婷婷六月久久综合丁香| 在线免费观看的www视频| 五月玫瑰六月丁香| 男女边吃奶边做爰视频| 狂野欧美激情性xxxx在线观看| 天天躁日日操中文字幕| 九九热线精品视视频播放| 国产一区二区亚洲精品在线观看| 99九九线精品视频在线观看视频| 青青草视频在线视频观看| 久久久久九九精品影院| 一级av片app| 蜜桃久久精品国产亚洲av| 日本三级黄在线观看| 午夜福利成人在线免费观看| 国产av一区在线观看免费| 国产日韩欧美在线精品| 成人午夜精彩视频在线观看| 中国美白少妇内射xxxbb| 国产精品久久久久久久电影| 国产精品,欧美在线| 日韩成人av中文字幕在线观看| av视频在线观看入口| 亚洲四区av| 精品无人区乱码1区二区| 欧美丝袜亚洲另类| 麻豆国产av国片精品| 午夜久久久久精精品| 日韩,欧美,国产一区二区三区 | 国产欧美日韩精品一区二区| 一个人看视频在线观看www免费| 亚洲中文字幕日韩| 男人舔女人下体高潮全视频| 国产精品久久久久久av不卡| 狠狠狠狠99中文字幕| 国模一区二区三区四区视频| 少妇猛男粗大的猛烈进出视频 | 熟女人妻精品中文字幕| 高清毛片免费观看视频网站| 国产69精品久久久久777片| 国产麻豆成人av免费视频| 简卡轻食公司| 亚洲人与动物交配视频| 欧美色视频一区免费| 性欧美人与动物交配| 成人性生交大片免费视频hd| 久久久a久久爽久久v久久| 久久精品国产亚洲av天美| 成人午夜精彩视频在线观看| 国产在线精品亚洲第一网站| 日韩欧美精品免费久久| 中国美女看黄片| 69av精品久久久久久| 51国产日韩欧美| 老司机影院成人| 中文字幕人妻熟人妻熟丝袜美| 老司机福利观看| 人人妻人人看人人澡| 美女被艹到高潮喷水动态| 日韩国内少妇激情av| 亚洲性久久影院| 波多野结衣高清无吗| 久久精品久久久久久噜噜老黄 | www.色视频.com| 日韩成人av中文字幕在线观看| 一进一出抽搐动态| 国产精品永久免费网站| 中文资源天堂在线| 99久久久亚洲精品蜜臀av| eeuss影院久久| 激情 狠狠 欧美| 亚洲最大成人手机在线| 最近最新中文字幕大全电影3| 小说图片视频综合网站| 免费av毛片视频| 国产黄片视频在线免费观看| 热99re8久久精品国产| 国产成人a区在线观看| 日本成人三级电影网站| 免费不卡的大黄色大毛片视频在线观看 | 激情 狠狠 欧美| 欧美色欧美亚洲另类二区| 人人妻人人看人人澡| 国产午夜福利久久久久久| 美女国产视频在线观看| 99国产精品一区二区蜜桃av| 直男gayav资源| 国产免费男女视频| 97人妻精品一区二区三区麻豆| 免费电影在线观看免费观看| 综合色丁香网| 久久久色成人| 精品99又大又爽又粗少妇毛片| 国产精品野战在线观看| 在线观看66精品国产| 只有这里有精品99| 亚洲熟妇中文字幕五十中出| 一个人观看的视频www高清免费观看| 可以在线观看的亚洲视频| 亚洲国产精品久久男人天堂| 看免费成人av毛片| 亚洲激情五月婷婷啪啪| 我的女老师完整版在线观看| 69人妻影院| 乱系列少妇在线播放| 国内久久婷婷六月综合欲色啪| 国产伦一二天堂av在线观看| 精华霜和精华液先用哪个| 最近视频中文字幕2019在线8| 熟女人妻精品中文字幕| 中文在线观看免费www的网站| 内射极品少妇av片p| 国产真实伦视频高清在线观看| a级毛片a级免费在线| 免费观看的影片在线观看| 蜜桃久久精品国产亚洲av| 91精品一卡2卡3卡4卡| 亚洲成人久久性| 波多野结衣高清作品| 亚洲五月天丁香| 啦啦啦啦在线视频资源| 成人综合一区亚洲| 白带黄色成豆腐渣| 亚洲欧美精品专区久久| 床上黄色一级片| 乱系列少妇在线播放| 中文在线观看免费www的网站| 亚洲在久久综合| 99国产极品粉嫩在线观看| 免费电影在线观看免费观看| 免费搜索国产男女视频| 秋霞在线观看毛片| 精品久久久久久久久久免费视频| 国产高清不卡午夜福利| 欧美bdsm另类| 中国国产av一级| 成人性生交大片免费视频hd| 看黄色毛片网站| 午夜福利高清视频| av免费在线看不卡| 亚洲aⅴ乱码一区二区在线播放| 91精品国产九色| 伦精品一区二区三区| 久久欧美精品欧美久久欧美| 国产精品久久久久久精品电影| 美女内射精品一级片tv| 国产精品蜜桃在线观看 | 熟女电影av网| 久久婷婷人人爽人人干人人爱| 最近手机中文字幕大全| 午夜久久久久精精品| 高清午夜精品一区二区三区 | 欧美极品一区二区三区四区| 国产午夜精品论理片| 国产亚洲5aaaaa淫片| 一个人免费在线观看电影| 桃色一区二区三区在线观看| 97超碰精品成人国产| 又粗又爽又猛毛片免费看| 欧美一区二区国产精品久久精品| 久久久久九九精品影院| 一本精品99久久精品77| 国内少妇人妻偷人精品xxx网站| 国产精品一区二区三区四区久久| 99久久精品国产国产毛片| 国产在视频线在精品| 熟女电影av网| 国产精品国产高清国产av| 日本五十路高清| 综合色av麻豆| 欧美xxxx性猛交bbbb| a级一级毛片免费在线观看| 日韩一本色道免费dvd| 国产成人福利小说| 国产精品一区www在线观看| 日韩大尺度精品在线看网址| 国产精品一区二区性色av| 国产精品人妻久久久影院| 欧美xxxx性猛交bbbb| 精品久久久久久久人妻蜜臀av| ponron亚洲| 男女边吃奶边做爰视频| 亚洲欧美中文字幕日韩二区| 欧美一级a爱片免费观看看| 国产极品精品免费视频能看的| 美女黄网站色视频| 日本一二三区视频观看| 最近最新中文字幕大全电影3| 亚洲中文字幕日韩| 精品午夜福利在线看| 国产精品美女特级片免费视频播放器| 国产精品电影一区二区三区| 中文字幕av成人在线电影| 狂野欧美激情性xxxx在线观看| 国产精品一二三区在线看| 欧美激情国产日韩精品一区| 村上凉子中文字幕在线| 日本免费a在线| 日韩一区二区三区影片| 级片在线观看| 欧美一级a爱片免费观看看| 插逼视频在线观看| 少妇被粗大猛烈的视频| 国产精品精品国产色婷婷| 99在线视频只有这里精品首页| 麻豆乱淫一区二区| 国产老妇伦熟女老妇高清| 色视频www国产| 一级黄色大片毛片| 久久久久久伊人网av| 非洲黑人性xxxx精品又粗又长| 欧美最黄视频在线播放免费| 成人毛片a级毛片在线播放| 综合色av麻豆| 国产精品久久久久久精品电影小说 | 国产高潮美女av| 久久精品影院6| 国产美女午夜福利| 少妇熟女aⅴ在线视频| 成人特级av手机在线观看| 黄色欧美视频在线观看| 菩萨蛮人人尽说江南好唐韦庄 | 亚洲精品色激情综合| 国产真实乱freesex| 亚洲av中文字字幕乱码综合| 热99在线观看视频| 99热6这里只有精品| 婷婷精品国产亚洲av| 久久精品国产亚洲av香蕉五月| 老司机影院成人| 不卡视频在线观看欧美| 国产精品一及| 能在线免费观看的黄片| 不卡视频在线观看欧美| 国产av不卡久久| 简卡轻食公司| 国产伦一二天堂av在线观看| 精品一区二区免费观看| 亚洲成人久久爱视频| 国产毛片a区久久久久| 国产av不卡久久| 国产精品蜜桃在线观看 | 3wmmmm亚洲av在线观看| 久久久色成人| 国产探花极品一区二区| 久久99热这里只有精品18| 国产精品久久久久久精品电影小说 | 日韩av不卡免费在线播放| 我要搜黄色片| 中文字幕人妻熟人妻熟丝袜美| 国产高潮美女av| 校园人妻丝袜中文字幕| 国产大屁股一区二区在线视频| 日本免费一区二区三区高清不卡| 搡老妇女老女人老熟妇| 又爽又黄a免费视频| 国产精品综合久久久久久久免费| or卡值多少钱| 久久草成人影院| 一个人看的www免费观看视频| 免费不卡的大黄色大毛片视频在线观看 | 久久久久久久亚洲中文字幕| 午夜福利高清视频| 女人十人毛片免费观看3o分钟| 亚洲在线自拍视频| 日韩人妻高清精品专区| 69人妻影院| 天堂网av新在线| 亚洲成人精品中文字幕电影| 欧美zozozo另类| 免费观看精品视频网站| 自拍偷自拍亚洲精品老妇| 国模一区二区三区四区视频| 午夜视频国产福利| 18禁黄网站禁片免费观看直播| 一个人观看的视频www高清免费观看| 免费人成在线观看视频色| 精品人妻视频免费看| 最近手机中文字幕大全| 人人妻人人看人人澡| 男女做爰动态图高潮gif福利片| 亚洲人与动物交配视频| 久久久久久久午夜电影| 美女内射精品一级片tv| 看黄色毛片网站| 99久久精品热视频| 一边摸一边抽搐一进一小说| 赤兔流量卡办理| 99久久精品热视频| 99热全是精品| 一级黄色大片毛片| 亚洲最大成人av| 精品久久久久久久久久久久久| 国产av在哪里看| 最后的刺客免费高清国语| 国产成人午夜福利电影在线观看| 国产精品av视频在线免费观看| 成人性生交大片免费视频hd| 亚洲色图av天堂| 欧美一区二区国产精品久久精品| 一本精品99久久精品77| 中文欧美无线码| 老司机影院成人| 成年版毛片免费区| 亚洲无线在线观看| 在线播放国产精品三级| 神马国产精品三级电影在线观看| 99久久精品国产国产毛片| 免费看光身美女| 长腿黑丝高跟| 九九爱精品视频在线观看| 亚洲精品乱码久久久久久按摩| 一区二区三区免费毛片| 国产精品不卡视频一区二区| 亚洲国产欧洲综合997久久,| 91精品国产九色| 日本免费一区二区三区高清不卡| 卡戴珊不雅视频在线播放| 欧美变态另类bdsm刘玥| 菩萨蛮人人尽说江南好唐韦庄 | 免费电影在线观看免费观看| 免费看日本二区| 天天躁夜夜躁狠狠久久av| 97热精品久久久久久| 伦精品一区二区三区| 国产国拍精品亚洲av在线观看| 成人亚洲精品av一区二区| 亚洲精品自拍成人| 亚洲精品国产av成人精品| 九草在线视频观看| 一个人免费在线观看电影| 日韩成人伦理影院| 91久久精品国产一区二区三区| 真实男女啪啪啪动态图| eeuss影院久久| 非洲黑人性xxxx精品又粗又长| 夜夜看夜夜爽夜夜摸| 69人妻影院| 日韩欧美在线乱码| 波多野结衣高清作品| 九九爱精品视频在线观看| 免费观看a级毛片全部| 精品人妻熟女av久视频| 国产熟女欧美一区二区| 欧美在线一区亚洲| 国产在线男女| 国产精品,欧美在线| 中文字幕免费在线视频6| 国产三级在线视频| 熟女人妻精品中文字幕| 午夜福利在线观看免费完整高清在 | 亚洲最大成人手机在线| 国产高清不卡午夜福利| 狠狠狠狠99中文字幕| 亚洲精华国产精华液的使用体验 | 亚洲第一电影网av| 哪个播放器可以免费观看大片| 久久精品夜色国产| 国产一级毛片七仙女欲春2| 性插视频无遮挡在线免费观看| 亚洲,欧美,日韩| 久久久久久久午夜电影| 天堂网av新在线| av免费在线看不卡| 免费黄网站久久成人精品| 日本色播在线视频| 国产视频内射| 日韩人妻高清精品专区| 欧美色视频一区免费| 你懂的网址亚洲精品在线观看 | av在线播放精品| 两性午夜刺激爽爽歪歪视频在线观看| 国产精品爽爽va在线观看网站| 成人一区二区视频在线观看| av又黄又爽大尺度在线免费看 | 亚洲高清免费不卡视频| 久久鲁丝午夜福利片| 日韩成人av中文字幕在线观看| 97超碰精品成人国产| 国产精品爽爽va在线观看网站| a级一级毛片免费在线观看| 欧美精品国产亚洲| 男女那种视频在线观看| 99热网站在线观看| 国产久久久一区二区三区| 亚洲精品日韩av片在线观看| 成年免费大片在线观看| 久久久午夜欧美精品| 日韩 亚洲 欧美在线| 边亲边吃奶的免费视频| 麻豆成人午夜福利视频| 国产一级毛片七仙女欲春2| 春色校园在线视频观看| 插逼视频在线观看| 亚洲婷婷狠狠爱综合网| 看十八女毛片水多多多| 免费人成视频x8x8入口观看| 成人高潮视频无遮挡免费网站| 免费观看人在逋| 国产伦一二天堂av在线观看| 听说在线观看完整版免费高清| 在线国产一区二区在线| 国国产精品蜜臀av免费| 一级毛片aaaaaa免费看小| 国产亚洲91精品色在线| 国产成人精品一,二区 | 99热这里只有精品一区| 亚洲精品成人久久久久久| 极品教师在线视频| 亚洲精品成人久久久久久| 国产一级毛片在线| 97人妻精品一区二区三区麻豆| 一级av片app| www日本黄色视频网| 精品日产1卡2卡| 3wmmmm亚洲av在线观看| 人妻久久中文字幕网| avwww免费| 亚洲图色成人| av又黄又爽大尺度在线免费看 | 插阴视频在线观看视频| 国产亚洲精品久久久久久毛片| 亚洲va在线va天堂va国产| 你懂的网址亚洲精品在线观看 | 99久久无色码亚洲精品果冻| 久久久久久久久久久免费av| 99久久无色码亚洲精品果冻| 国语自产精品视频在线第100页| 毛片一级片免费看久久久久| 免费无遮挡裸体视频| 亚洲五月天丁香| 三级男女做爰猛烈吃奶摸视频| 亚洲精品影视一区二区三区av| 夫妻性生交免费视频一级片| 18禁在线无遮挡免费观看视频| 女同久久另类99精品国产91| 日本一二三区视频观看| 久久精品91蜜桃| 好男人在线观看高清免费视频| 欧美一区二区精品小视频在线| 热99re8久久精品国产| 亚洲精品自拍成人| 夜夜夜夜夜久久久久| 精华霜和精华液先用哪个| 久久精品国产亚洲av天美|