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

    基于GPU的密碼S盒代數(shù)性質(zhì)評(píng)估方法

    2022-09-25 08:42:32蔡婧雯韋永壯劉爭(zhēng)紅
    計(jì)算機(jī)應(yīng)用 2022年9期
    關(guān)鍵詞:均勻度內(nèi)核線程

    蔡婧雯,韋永壯*,劉爭(zhēng)紅

    (1.廣西密碼學(xué)與信息安全重點(diǎn)實(shí)驗(yàn)室(桂林電子科技大學(xué)),廣西桂林 541004;2.廣西無(wú)線寬帶通信與信號(hào)處理重點(diǎn)實(shí)驗(yàn)室(桂林電子科技大學(xué)),廣西桂林 541004)

    0 引言

    密碼S 盒作為對(duì)稱密碼算法的核心部件,主要提供了必要的非線性變換,其代數(shù)性質(zhì)往往決定著密碼算法的安全強(qiáng)度。伴隨著超級(jí)計(jì)算機(jī)其計(jì)算能力的迅速提升,特別是抵抗未來(lái)的量子計(jì)算攻擊[1],高強(qiáng)度密碼算法設(shè)計(jì)中對(duì)S 盒的輸入及輸出規(guī)模提出了新的要求,比如基于非線性反饋移位寄存 器(Nonlinear Feedback Shift Register,NFSR)或ARX(Addition-Rotation-XOR)操作部件等方法構(gòu)造出16 比特或32 比特的大狀態(tài)密碼S 盒。美國(guó)國(guó)家標(biāo)準(zhǔn)與技術(shù)研究院(National Institute of Standards and Technology,NIST)發(fā)起輕量級(jí)密碼算法公開征集[2],最終入圍算法中SPARKLE[3]及GIFT-COFB(COmbined FeedBack)[4]等算法均采用了32 比特或者64 比特大狀態(tài)密碼S 盒。同年,在中國(guó)密碼學(xué)會(huì)舉辦的全國(guó)密碼算法設(shè)計(jì)競(jìng)賽[5]中,徐洪等[6]基于16 級(jí)NFSR 迭代構(gòu)造了16 比特S 盒;田甜等[7]基于NFSR 設(shè)計(jì)了32 比特S 盒;2020 年美洲密碼會(huì)上,Beierle 等[8]基于ARX 結(jié)構(gòu)構(gòu)造了64比特S 盒。注意到,密碼S 盒的安全性與其安全性指標(biāo)息息相關(guān),傳統(tǒng)安全性指標(biāo)包括差分均勻度[9]、非線性度[10-11]、透明階(Revised Transparency Order,RTO)[12]、飛來(lái)器連接表[13](Boomerang Connectivity Table,BCT)等。這些指標(biāo)與相應(yīng)的密碼攻擊密切相關(guān),如差分均勻度、非線性度和透明階分別刻畫了S 盒抵抗差分密碼分析[14]、線性密碼分析[15]及差分功耗攻擊(Differential Power Attack,DPA)[16]的能力。另一方面,對(duì)于n比特輸入及n比特輸出的密碼S 盒,當(dāng)n比較大時(shí)(如n>15 時(shí))評(píng)估S 盒的各個(gè)安全性指標(biāo)則較為困難,比如傳統(tǒng)求解密碼S 盒的差分均勻度、非線性度及透明階時(shí)間復(fù)雜度分別約為O(23n)、O(23n)及O(23n·n2)。這些求解因搜索空間大,從而導(dǎo)致花銷時(shí)間太長(zhǎng)等問(wèn)題。如何快速評(píng)估密碼S 盒的代數(shù)性質(zhì)是目前研究的熱點(diǎn)之一。

    為了解決計(jì)算資源瓶頸,圖形處理器(Graphics Processing Unit,GPU)應(yīng)運(yùn)而生。GPU 主要應(yīng)用于圖像處理、視頻音頻處理、計(jì)算生物學(xué)等領(lǐng)域上。而利用GPU 解決密碼學(xué)問(wèn)題最早工作是Kedem 等[17]使用PixelFlow 圖像引擎快速破解了UNIX 系統(tǒng)密鑰;Manavski[18]利用計(jì)算統(tǒng)一設(shè)備架構(gòu)(Compute Unified Device Architecture,CUDA)進(jìn)行高級(jí)加密標(biāo)準(zhǔn)(Advanced Encryption Standard,AES)加速;Cheong等[19]在具有Kepler 架構(gòu)上提出了加速分組密碼算法國(guó)際數(shù)據(jù)加密算法(International Data Encryption Algorithm,IDEA),進(jìn)一步提高加密吞吐量;Yeoh 等[20]提出了一種基于GPU 的分支定界算法。如何基于GPU 快速評(píng)估密碼S 盒的安全強(qiáng)度仍有待進(jìn)一步研究。

    本文基于CPU-GPU 異構(gòu)結(jié)構(gòu),對(duì)密碼S 盒的差分均勻度、非線性度及透明階提出求解優(yōu)化方法,實(shí)現(xiàn)多線程并行計(jì)算,提出一種快速求解差分均勻度、非線性度及透明階的方法。測(cè)試結(jié)果表明,與基于中央處理器(Central Processing Unit,CPU)實(shí)現(xiàn)相比,基于CPU-GPU 異構(gòu)結(jié)構(gòu)實(shí)現(xiàn)效率得到大幅度提升。本文方法利用單塊GPU 分別計(jì)算差分均勻度、非線性度及透明階所花銷的時(shí)間與傳統(tǒng)方法相比節(jié)省了90.28%、78.57%、60%。

    1 預(yù)備知識(shí)

    定義1設(shè)n比特輸入、m比特輸出的S 盒記為S(x)=(f1(x),f2(x),…,fm(x)):→,其中fi(x) 為→F2的布爾函數(shù),記為密碼S 盒的第i個(gè)分量函數(shù)(i=1,2,…,m)。本 文考慮n=m=16 的16 比特S 盒。

    定義2差分均勻度[9]。設(shè)n比特輸入、n比特輸出的S盒記為S,對(duì)任意的輸入差分α∈和輸出差分β∈,其中差分對(duì)解的個(gè)數(shù)為:

    則差分均勻度定義為:

    當(dāng)差分均勻度越小,S 盒的差分分布更均勻,安全性更好。

    定義3非線性度[10-11]。一個(gè)n×n的S 盒的非線性度為S 盒的所有分量函數(shù)的非零線性組合中最小的非線性度,即:

    定義4透明階(RTO)[12]。對(duì)于任意n比特輸入、n比特輸出的S 盒,該S 盒的透明階記為:

    當(dāng)透明階越小時(shí),抵抗差分功耗攻擊的能力越強(qiáng),安全性越好。

    定義5S 盒的差分概率[21]。對(duì)于一個(gè)n比特輸入、n比特輸出的密碼S 盒,對(duì)于輸入差分α∈和輸出差分β∈,存在差分概率

    則稱輸入差分α經(jīng)過(guò)S 盒后將以概率PS(α→β)得到輸出差分β。

    定義6S 盒的線性概率[21]。定義一個(gè)S 盒S:→,定義NS(θ,λ)=#{x∈:θ·x=λ·S(x)}構(gòu)造密碼S 盒的線性逼近表,其中,固定輸入掩碼θ∈,得到輸出掩碼λ的概率為:

    即固定輸入掩碼θ,輸出掩碼λ,隨機(jī)給定輸入x,則θ·x=λ·S(x)以概率PLS(θ→λ)成立得到。

    2 CUDA的并行計(jì)算

    2.1 GPU與CUDA

    隨著人工智能、大數(shù)據(jù)等計(jì)算領(lǐng)域的不斷發(fā)展,計(jì)算復(fù)雜度越來(lái)越大,圖像處理器(GPU)從幫助CPU 做圖像和圖形運(yùn)算轉(zhuǎn)至海量數(shù)據(jù)處理上,涉及云計(jì)算、生物計(jì)算、天文學(xué)等多個(gè)領(lǐng)域,逐步成為了計(jì)算領(lǐng)域的研究熱點(diǎn)。

    由于CPU 與GPU 所應(yīng)用的場(chǎng)景不同,兩者的架構(gòu)大不相同。從圖1 中CPU 與GPU 的架構(gòu)相比可知,CPU 與GPU 有如下區(qū)別:

    圖1 CPU與GPU的架構(gòu)區(qū)別Fig.1 Difference between CPU and GPU architectures

    1)GPU 采用了若干個(gè)算術(shù)邏輯單元(Arithmetic and Logic Unit,ALU)和超長(zhǎng)的流水線,可以同時(shí)處理多個(gè)線程;然而CPU 擁有少量但很強(qiáng)大的算術(shù)邏輯單元,可以在很少的時(shí)間周期內(nèi)完成運(yùn)算。

    2)CPU 具有強(qiáng)大的控制邏輯單元,在運(yùn)算過(guò)程中提供邏輯預(yù)測(cè)能力降低延時(shí);而GPU 控制能力稍遜色于CPU,運(yùn)算過(guò)程中可以將多個(gè)訪問(wèn)合并成較少的訪問(wèn)。

    3)CPU 有大量的緩存空間降低計(jì)算延時(shí),而GPU 只有少量的緩存空間,與CPU 的緩存空間功能不同,GPU 的緩存空間將線程所需要訪問(wèn)的相同數(shù)據(jù)合并訪問(wèn)動(dòng)態(tài)隨機(jī)存取存儲(chǔ)器。

    從以上而可知,CPU 適合于需要較大的緩存空間且復(fù)雜控制邏輯的通信密集型運(yùn)算;GPU 則適合于邏輯分支簡(jiǎn)單、計(jì)算量大的計(jì)算密集型運(yùn)算。另外,從圖2 可以得知,在計(jì)算密集型任務(wù)時(shí),GPU 所需耗時(shí)占CPU 所需耗時(shí)的22%;而在計(jì)算通信密集的任務(wù)時(shí),CPU 計(jì)算所需耗時(shí)較少,約為GPU 計(jì)算所需耗時(shí)的30%。在一些具有計(jì)算密集要求且邏輯分支較為簡(jiǎn)單的計(jì)算任務(wù)中,GPU 的處理能力比CPU 具有更大的優(yōu)勢(shì)。

    圖2 CPU與GPU任務(wù)耗時(shí)比較Fig.2 Task time consumption comparison of CPU and GPU

    GPU 的造價(jià)和功耗與相同計(jì)算能力的CPU 相比,GPU 的造價(jià)和功耗相對(duì)較低。在計(jì)算領(lǐng)域中構(gòu)建CPU 集群的超級(jí)計(jì)算機(jī),造價(jià)昂貴。根據(jù)摩爾定律可以得知當(dāng)CPU 計(jì)算速度達(dá)到一定程度時(shí)提升空間受限,GPU 的出現(xiàn)滿足了需要計(jì)算大量數(shù)據(jù)而無(wú)法使用巨型計(jì)算機(jī)的用戶需求。

    目前應(yīng)用較為廣泛的GPU 并行編程平臺(tái)有CUDA、OpenCL 等。2006 年NVIDIA 公司推出并行開發(fā)平臺(tái)CUDA,支持C、C++、Java、Python 等多種主流編程語(yǔ)言,便于各領(lǐng)域進(jìn)行并行計(jì)算操作。CUDA 使用了具有很強(qiáng)的并行計(jì)算特點(diǎn)的單指令多線程(Single Instruction Multiple Thread,SIMT)的執(zhí)行模型,模型在執(zhí)行過(guò)程中,構(gòu)建CPU 與GPU 異構(gòu)架構(gòu),其中:CPU 主要負(fù)責(zé)串行計(jì)算工作,完成較為復(fù)雜的邏輯控制及通信密集的運(yùn)算;GPU 主要負(fù)責(zé)并行計(jì)算工作,完成運(yùn)算量大且計(jì)算任務(wù)較為簡(jiǎn)單的計(jì)算密集型工作。經(jīng)過(guò)測(cè)試及多方考量,本文選用CUDA 作為本文的并行開發(fā)平臺(tái),構(gòu)造CPU-GPU 異構(gòu)架構(gòu)進(jìn)行測(cè)試。

    2.2 線程塊的分配

    為了發(fā)揮GPU 的最大并行計(jì)算效率,在執(zhí)行內(nèi)核函數(shù)過(guò)程中,需要合理配置線程塊數(shù)量及每塊線程塊中的線程數(shù)量。線程塊的數(shù)目由配置參數(shù)所劃分的網(wǎng)格所決定,通常為32 的倍數(shù)最佳。通過(guò)測(cè)試可知,將線程塊的線程數(shù)量為512時(shí)具有最大的計(jì)算能力。

    2.3 并行計(jì)算的影響因素

    總結(jié)前文所述的GPU 與CUDA 的計(jì)算特點(diǎn),影響并行計(jì)算效率的因素可以總結(jié)為以下3 點(diǎn):

    1)減少CPU 與GPU 之間的數(shù)據(jù)傳輸。由于CPU 與GPU使用不同的內(nèi)存空間,CPU 與GPU 在數(shù)據(jù)交換過(guò)程中需要通過(guò)計(jì)算機(jī)總線,造成了額外的時(shí)間花銷,因此在利用GPU 進(jìn)行并行計(jì)算時(shí),應(yīng)盡量避免減少CPU 與GPU 之間的數(shù)據(jù)傳輸。

    2)減少訪問(wèn)GPU 全局內(nèi)存。為了減少內(nèi)存訪問(wèn)產(chǎn)生的時(shí)延和消耗,在GPU 中應(yīng)盡量減少過(guò)多的跳躍式訪問(wèn),最大限度減少因?qū)PU 內(nèi)存訪問(wèn)而造成的延遲。

    3)合理的資源配置。為了提高并行計(jì)算的效率,合理設(shè)置線程塊內(nèi)的線程數(shù)量,最大限度地利用線程處理計(jì)算任務(wù),減少資源的浪費(fèi)。

    3 NBC算法

    NBC 算法為中國(guó)密碼學(xué)會(huì)舉辦的全國(guó)密碼算法設(shè)計(jì)競(jìng)賽分組算法[11]第二輪入選算法之一,其采用廣義Feistel 結(jié)構(gòu)[12],算法加密共有三種模式,具體如表1 所示。

    表1 NBC算法的三種模式Tab.1 Three modes of NBC algorithm

    本文使用的算法是數(shù)據(jù)分組長(zhǎng)度為128 比特、密鑰長(zhǎng)度為128 比特的NBC 算法。設(shè)第i輪的輸入為Xi=輸出為NBC-128/128 算法結(jié)構(gòu)如圖3 所示。

    圖3 八分支的1輪NBC-128/128結(jié)構(gòu)Fig.3 One-round NBC-128/128 structure with 8 branches

    NBC-128/128 算法的S 盒采用16 級(jí)NFSR來(lái)構(gòu)造16 比特S 盒,S 盒構(gòu)造圖如圖4 所示。設(shè)S 盒的16 比特的輸入為S0S1…S15,當(dāng)全體內(nèi)部狀態(tài)經(jīng)過(guò)迭代20 輪后形成S 盒輸出。

    圖4 NBC-128算法的S盒構(gòu)造Fig.4 S-box structure of NBC-128 algorithm

    算法設(shè)計(jì)者稱構(gòu)造出來(lái)的S 盒的差分均勻度Diff(S)=22,非線性度NL(S)=31 982,透明階RTO=15.982 6。

    4 基于GPU的16比特密碼S盒代數(shù)性質(zhì)評(píng)估

    由于在CPU 下求解差分均勻度、非線性度及透明階的算法效率較低,在本章中,將傳統(tǒng)求解密碼S 盒代數(shù)性質(zhì)評(píng)估方法進(jìn)行優(yōu)化,分別討論基于單GPU 模式和多GPU 模式下將內(nèi)核函數(shù)切片至多線程中,實(shí)現(xiàn)多線程并行化計(jì)算。

    4.1 單GPU對(duì)16比特密碼S盒性質(zhì)評(píng)估

    根據(jù)共享式內(nèi)存的結(jié)構(gòu)特點(diǎn)和對(duì)S 盒性質(zhì)評(píng)估的求解流程,本文提出了單塊GPU 環(huán)境下的CPU-GPU 異構(gòu)模式,并行架構(gòu)如圖5 所示。

    圖5 CPU-GPU異構(gòu)并行流程Fig.5 CPU-GPU heterogeneous parallel flowchart

    程序在運(yùn)行時(shí)控制一塊GPU,創(chuàng)建多個(gè)線程共同完成計(jì)算任務(wù)。具體步驟如下所示:

    1)檢測(cè)顯卡設(shè)備。函數(shù)cudaSetDevice()表示檢測(cè)主機(jī)設(shè)備的顯卡個(gè)數(shù),當(dāng)檢測(cè)到主機(jī)存在可使用的顯卡時(shí),將對(duì)算法進(jìn)行CUDA 并行計(jì)算做好準(zhǔn)備;

    2)讀取數(shù)據(jù)并復(fù)制入GPU。采用cudaMalloc()函數(shù)在設(shè)備Device 中開辟計(jì)算中所需要參數(shù)的空間。由于GPU 在計(jì)算過(guò)程中,無(wú)法直接讀取CPU 內(nèi)存中的數(shù)據(jù),故在計(jì)算前需要在設(shè)備Decive 開辟相應(yīng)的空間。

    3)當(dāng)Device 中開辟了相應(yīng)的空間大小后,使用cudaMemcpy()函數(shù)將所需要的參數(shù)S 盒復(fù)制進(jìn)入GPU 內(nèi)。

    4)計(jì)算內(nèi)核函數(shù)。偽代碼中存在3 個(gè)不同的內(nèi)核函數(shù),分別為differenceUniformity()、degreeOfNolinearity()及calculateRTO(),其中:differenceUniformity()為計(jì)算差分均勻度的內(nèi)核函數(shù);degreeOfNolinearity()為計(jì)算非線性度的內(nèi)核函數(shù);calculateRTO()為計(jì)算透明階的內(nèi)核函數(shù)。

    5)在內(nèi)核函數(shù)中,<<<Block,Thread>>>表示在啟動(dòng)內(nèi)核函數(shù)時(shí),分配Block個(gè)線程組,每個(gè)線程組中分配Thread個(gè)線程,故共分配Block*Thread線程總數(shù)。通過(guò)合理設(shè)置線程組和線程數(shù)量,才能更好地發(fā)揮GPU 的計(jì)算能力。本文使用的是一個(gè)線程處理一個(gè)分組,例如當(dāng)處理100 組數(shù)據(jù)時(shí),需要在GPU 內(nèi)分配100 個(gè)線程,故本文計(jì)算16 比特S 盒的密碼性質(zhì)中,共需要處理65 536 個(gè)分組數(shù)據(jù),使用了128 個(gè)線程塊,其中每個(gè)線程塊512 個(gè)線程。

    6)檢查并返回結(jié)果。當(dāng)每一個(gè)線程完成了內(nèi)核函數(shù)中的計(jì)算任務(wù)時(shí),使用cudaGetLastError()函數(shù)檢查內(nèi)核函數(shù)在計(jì)算過(guò)程中是否存在錯(cuò)誤:若存在錯(cuò)誤,將錯(cuò)誤返回至CPU中;若不存在,利用函數(shù)cudaMemcpy()將計(jì)算結(jié)果返回至CPU 中,計(jì)算結(jié)束。

    求解復(fù)雜度分析如下:

    1)由差分均勻度的定義可以得知:針對(duì)n比特輸入、n比特輸出的密碼S 盒,傳統(tǒng)求解差分均勻度需要遍歷輸入差分α∈、輸出差分β∈及x∈三個(gè)變量,時(shí)間復(fù)雜度約為O(23n)。根據(jù)GPU 并行計(jì)算的特性,使用切片技術(shù)對(duì)求解差分均勻度的最外層循環(huán)分解到各個(gè)線程中并行,即除最外層循環(huán)外部分設(shè)為內(nèi)核函數(shù),此時(shí)求解差分均勻度的時(shí)間復(fù)雜度降低至O(22n)。為了進(jìn)一步提高效率,減少計(jì)算邏輯分支數(shù),再將遍歷的輸出差分β循環(huán)放置內(nèi)核函數(shù)外,此在GPU 內(nèi)計(jì)算的內(nèi)核函數(shù)的時(shí)間復(fù)雜度將降低至O(2n)。

    2)對(duì)求解非線性度及透明階進(jìn)行求解分析。傳統(tǒng)求解非線性度及透明階的時(shí)間復(fù)雜度為O(23n)、O(23n·n2)。利用相同的切片技術(shù),將求解最外層循環(huán)分解到各個(gè)線程中,求解過(guò)程中利用線程索引對(duì)應(yīng)最外層循環(huán)所遍歷的值,此時(shí)求解非線性度及透明階的時(shí)間復(fù)雜度降低至O(22n)、O(22n·n2)。另外再將一層循環(huán)放在內(nèi)核函數(shù)外,最終GPU 內(nèi)計(jì)算非線性度及透明階的內(nèi)核函數(shù)的時(shí)間復(fù)雜度將降低至O(2n)、O(2n·n2),與傳統(tǒng)求解方法相比,該方法的時(shí)間復(fù)雜度降低了兩個(gè)指數(shù)級(jí),節(jié)省了求解時(shí)間花銷。

    算法1 測(cè)試主程序。

    輸入 S 盒;

    輸出 差分均勻度,非線性度,透明階。

    4.2 多GPU對(duì)大狀態(tài)S盒性質(zhì)評(píng)估

    4.1 節(jié)分析了在CPU-GPU 異構(gòu)計(jì)算結(jié)構(gòu)下,對(duì)16 比特S盒安全性指標(biāo)測(cè)評(píng)比在傳統(tǒng)CPU 計(jì)算下所具有的時(shí)間優(yōu)勢(shì),在相同的實(shí)驗(yàn)條件下,使用單塊GPU 構(gòu)建的CPU-GPU 異構(gòu)計(jì)算比傳統(tǒng)CPU 計(jì)算時(shí)間節(jié)省90.28%。但對(duì)于n比特輸入、n比特輸出的密碼S 盒,當(dāng)n比較大時(shí)(如n>15 時(shí)),由于計(jì)算搜索空間大,運(yùn)算量大,單GPU 計(jì)算時(shí)間仍然較長(zhǎng),故提出在多GPU 環(huán)境下,對(duì)評(píng)估NBC 算法的16 比特S 盒的差分均勻度、非線性度等安全性指標(biāo)方案并行化研究,對(duì)計(jì)算過(guò)程中涉及的數(shù)據(jù)傳輸過(guò)程進(jìn)行研究與優(yōu)化。分析并行化計(jì)算所遇到的瓶頸主要在數(shù)據(jù)傳輸過(guò)程,在結(jié)果保證正確性的基礎(chǔ)上,調(diào)整程序的傳輸方式,由同步傳輸調(diào)整至異步傳輸,且利用多流技術(shù)與異步傳輸相結(jié)合逐步提高加速比。最后通過(guò)實(shí)現(xiàn)分析說(shuō)明基于多GPU 架構(gòu)下對(duì)大狀態(tài)S 盒的安全性指標(biāo)計(jì)算性能。

    在使用多GPU 構(gòu)架中,選擇單個(gè)節(jié)點(diǎn)連接到高速串行計(jì)算機(jī)擴(kuò)展總線標(biāo)準(zhǔn)(Peripheral Component Interconnect express,PCIe)總線上,具體架構(gòu)如圖6 所示。程序在運(yùn)行時(shí)使用函數(shù)cudaSetDecice()對(duì)GPU 設(shè)備組上的各設(shè)備進(jìn)行綁定,使得每個(gè)線程管理一個(gè)GPU,實(shí)現(xiàn)多個(gè)GPU 并行工作。與單GPU 結(jié)構(gòu)相比,多GPU 結(jié)構(gòu)可以開辟更多的線程,運(yùn)算速度得到進(jìn)一步提升。

    圖6 多GPU節(jié)點(diǎn)架構(gòu)Fig.6 Multi-GPU node architecture

    由于同步傳輸?shù)牟⑿谢?jì)算中,傳輸數(shù)據(jù)占用了大量的時(shí)間。本節(jié)利用多流技術(shù)與異步技術(shù)相結(jié)合,在計(jì)算過(guò)程中使計(jì)算過(guò)程與數(shù)據(jù)傳輸兩個(gè)步驟進(jìn)行重疊,從而減少一部分時(shí)間的開銷。有無(wú)重疊優(yōu)化的時(shí)間開銷對(duì)比如圖7 所示。在無(wú)重疊優(yōu)化時(shí),由于默認(rèn)只有一個(gè)流隊(duì)列,此時(shí)所有的計(jì)算過(guò)程皆為串行執(zhí)行。先對(duì)數(shù)據(jù)傳輸至GPU 的全局內(nèi)存內(nèi),傳輸完畢后再進(jìn)行數(shù)據(jù)計(jì)算,等待GPU 內(nèi)的所有線程計(jì)算完畢后再將結(jié)果復(fù)制回CPU 內(nèi)。作為對(duì)比,在使用重疊技術(shù)進(jìn)行時(shí)間優(yōu)化后,當(dāng)一個(gè)流隊(duì)列在計(jì)算部分?jǐn)?shù)據(jù)的同時(shí),另一個(gè)流隊(duì)列可以對(duì)剩下數(shù)據(jù)進(jìn)行傳輸至GPU 內(nèi)等待計(jì)算。當(dāng)一個(gè)流上的數(shù)據(jù)計(jì)算完畢后,利用另一個(gè)流隊(duì)列傳輸回CPU 內(nèi),下一個(gè)流隊(duì)列等待數(shù)據(jù)傳輸。計(jì)算與傳輸時(shí)間重疊技術(shù)的優(yōu)化既能保持計(jì)算任務(wù)仍按照串行執(zhí)行,又能掩蓋GPU 與CPU 數(shù)據(jù)傳輸之間所帶來(lái)的大量時(shí)間開銷,從而進(jìn)一步減少程序所需要的執(zhí)行時(shí)間,提高并行效率。

    圖7 有無(wú)重疊優(yōu)化的時(shí)間開銷對(duì)比Fig.7 Comparison of time cost with and without overlapping optimization

    利用多GPU 對(duì)大狀態(tài)S 盒進(jìn)行評(píng)估過(guò)程具體如下:

    1)在CPU 端獲取已有的GPU 設(shè)備數(shù)量和每個(gè)GPU 設(shè)備信息。利用CUDA 中自帶的函數(shù)cudaGetDeviceCount(&ngpus),讀取已有的GPU 設(shè)備數(shù)量,并將GPU 數(shù)量信息存儲(chǔ)在變量ngpus中,可通過(guò)設(shè)備號(hào)dev進(jìn)行選擇使用GPU設(shè)備。

    2)在同一節(jié)點(diǎn)上的GPU 設(shè)備構(gòu)成GPU 設(shè)備組,GPU 設(shè)備組內(nèi)的GPU 設(shè)備直接進(jìn)行通信和數(shù)據(jù)傳輸。

    3)在CPU 端進(jìn)一步準(zhǔn)備計(jì)算所需要的數(shù)據(jù)集,根據(jù)GPU設(shè)備組的數(shù)量,將數(shù)據(jù)平分至各GPU 設(shè)備上,另外在CPU 端設(shè)置S 盒,在CPU 端將S 盒以結(jié)構(gòu)體的形式傳輸至GPU 設(shè)備組中對(duì)應(yīng)的常量存儲(chǔ)區(qū)中,S 盒參數(shù)都在對(duì)應(yīng)的GPU 設(shè)備運(yùn)行過(guò)程中將會(huì)被核函數(shù)多次調(diào)用。

    4)在CPU 端設(shè)置循環(huán)遍歷所有的GPU 設(shè)備,將GPU 設(shè)備組中的GPU 分別置于對(duì)應(yīng)并行流上,通過(guò)對(duì)工作流在不同時(shí)間下的操作和阻塞實(shí)現(xiàn)GPU 設(shè)備的異步,設(shè)置CUDA 工作流Steam 的異步操作隱藏了部分訪問(wèn)延遲和實(shí)現(xiàn)了任務(wù)的并發(fā)執(zhí)行,減少了數(shù)據(jù)處理時(shí)間。

    5)明確每個(gè)核函數(shù)分配的變量和變量空間,利用函數(shù)cudaMemory()將數(shù)據(jù)以異步方式傳輸至GPU 設(shè)備組中對(duì)應(yīng)的GPU 上。

    6)為了確保核函數(shù)運(yùn)行時(shí)有較好的性能,延用上一節(jié)對(duì)GPU 的線程數(shù)分配,使用了128 個(gè)線程組,其中每個(gè)線程組512 個(gè)線程,共計(jì)65 536 個(gè)線程數(shù)。

    7)核函數(shù)完成線程配置后,數(shù)據(jù)根據(jù)“分而治之”思想,將輸入數(shù)據(jù)劃分成多個(gè)子集分別復(fù)制。由于每個(gè)問(wèn)題都是獨(dú)立的,所以分別安排在不同的并行流中進(jìn)行計(jì)算,不同的流之間輸出傳輸于另一個(gè)流的核計(jì)算進(jìn)行重疊。

    8)當(dāng)線程中循環(huán)遍歷完所有的塊,完成內(nèi)核函數(shù)Kernel的計(jì)算后,利用重疊流的思想保證每個(gè)線程計(jì)算后優(yōu)先傳輸至CPU 內(nèi)。

    9)當(dāng)所有線程都計(jì)算完畢后,CPU 端對(duì)GPU 設(shè)備組的各GPU 設(shè)備返回的結(jié)果進(jìn)行統(tǒng)一歸總,并按照規(guī)定的格式進(jìn)行輸出。

    算法2 多GPU 測(cè)試主程序。

    輸入 S 盒;

    輸出 差分均勻度,非線性度,透明階。

    5 測(cè)試與結(jié)果分析

    5.1 測(cè)試環(huán)境

    本文實(shí)驗(yàn)環(huán)境所使用的CPU 為Intel Xeon Silver 4210 2.20 GHz;GPU 為NVIDIA Quadro RTX 8000;在多GPU 環(huán)境下,共使用4 塊相同型號(hào)的GPU,且顯卡型號(hào)為NVIDIA Quadro RTX 8000;操作系統(tǒng)為Ubuntu 18.04.4 LTS,64 bits;編程環(huán)境為CUDA 7.0、GCC 7.5.0。本實(shí)驗(yàn)的CPU 代碼用的C 語(yǔ)言進(jìn)行編寫,GPU 代碼用CUDA C 進(jìn)行編寫。

    5.2 測(cè)試結(jié)果

    本次測(cè)試是針對(duì)NBC-128/128 算法的16 比特S 盒分別進(jìn)行差分均勻度、非線性度和透明階運(yùn)算,其中測(cè)試可得NBC 算法的差分均勻度為Diff(S)=22,非線性度為NL=31 982,透明階RTO=15.982 6,運(yùn)行時(shí)間如圖8 所示。

    圖8 對(duì)比CPU、單塊GPU及多塊GPU下的運(yùn)行時(shí)間Fig.8 Comparison of running time under CPU,single GPU and multi-GPU

    通過(guò)以上實(shí)驗(yàn)結(jié)果表明,在相同的實(shí)驗(yàn)條件下,使用GPU 測(cè)試16 比特S 盒差分均勻度所用時(shí)間比在CPU 測(cè)試16比特S 盒所用時(shí)間約減少90.28%;使用GPU 測(cè)試16 比特S盒的非線性度所用時(shí)間比在CPU 測(cè)試16 比特S 盒所用時(shí)間約減少78.57%;使用GPU 測(cè)試16 比特S 盒透明階所用時(shí)間比在CPU 測(cè)試16 比特S 盒所用時(shí)間約減少60%。實(shí)驗(yàn)結(jié)果證明使用GPU 測(cè)試大比特S 盒性質(zhì)所消耗時(shí)間明顯少于使用CPU 測(cè)試大比特S 盒性質(zhì)所用的時(shí)間。在使用多GPU 并行計(jì)算的架構(gòu)下,在相同實(shí)驗(yàn)條件下,使用多GPU 測(cè)試差分均勻度所用時(shí)間比單GPU 測(cè)試所用時(shí)間約減少99.52%;使用多GPU 測(cè)試非線性度所用時(shí)間比單GPU 測(cè)試所用時(shí)間約減少91.67%;使用多GPU 測(cè)試透明階所用時(shí)間比單GPU 測(cè)試所用時(shí)間約減少78.13%,使用多塊GPU 并行計(jì)算的計(jì)算速率明顯高于單塊GPU 計(jì)算速率。

    通過(guò)密碼S 盒的差分概率定義可知,當(dāng)輸入尺寸n比較大時(shí)(如n>15 時(shí)),需要遍歷輸入差分α∈、輸出差分β∈及x∈三個(gè)變量,所以求解差分概率所需的時(shí)間復(fù)雜度約為O(23n)。類似地,當(dāng)求解線性概率時(shí),同樣需要遍歷3 個(gè)變量,分別是輸入掩碼θ、輸出掩碼λ及輸入x,即時(shí)間復(fù)雜度也約為O(23n)。注意到,利用切片技術(shù)對(duì)差分概率及線性概率的計(jì)算過(guò)程可以分解到各個(gè)線程中并行計(jì)算。因而,求解差分概率及線性概率與求解差分均勻度方法類似,預(yù)計(jì)所花銷的時(shí)間大致相當(dāng),限于篇幅,本文不再贅述。

    6 結(jié)語(yǔ)

    本文基于CPU-GPU 結(jié)構(gòu),結(jié)合差分均勻度、非線性度等計(jì)算特征,將內(nèi)核函數(shù)利用切片技術(shù)拆分至多線程上,實(shí)現(xiàn)多線程并行計(jì)算,并由此提出快速評(píng)估密碼S 盒代數(shù)性質(zhì)新方法。在單塊GPU 及4 塊GPU 環(huán)境下對(duì)NBC-128/128 密碼算法的S 盒進(jìn)行差分均勻度、非線性度及透明階3 個(gè)性質(zhì)計(jì)算,實(shí)驗(yàn)結(jié)果證實(shí):與基于CPU 的實(shí)現(xiàn)環(huán)境相比,基于單塊GPU 所構(gòu)建的CPU-GPU 架構(gòu)的實(shí)現(xiàn)效率得到了顯著的提升,即計(jì)算差分均勻度、非線性度及透明階分別節(jié)省了90.28%、78.57%、60%的時(shí)間。下一步的研究工作可以考慮針對(duì)32 比特、64 比特等大狀態(tài)的密碼S 盒,基于CPU-GPU 結(jié)構(gòu)進(jìn)行安全性評(píng)估。

    猜你喜歡
    均勻度內(nèi)核線程
    低播量下雜交稻產(chǎn)量形成對(duì)種植均勻度的響應(yīng)
    作物研究(2023年2期)2023-05-28 13:44:14
    萬(wàn)物皆可IP的時(shí)代,我們當(dāng)夯實(shí)的IP內(nèi)核是什么?
    強(qiáng)化『高新』內(nèi)核 打造農(nóng)業(yè)『硅谷』
    均勻度控制不佳可致肉種雞晚產(chǎn)
    基于嵌入式Linux內(nèi)核的自恢復(fù)設(shè)計(jì)
    Linux內(nèi)核mmap保護(hù)機(jī)制研究
    淺談linux多線程協(xié)作
    錦綸長(zhǎng)絲染色均勻度判色新方法
    復(fù)方丹參片中冰片的含量均勻度研究
    中成藥(2014年10期)2014-02-28 22:29:24
    Linux線程實(shí)現(xiàn)技術(shù)研究
    av天堂中文字幕网| 国产在线男女| 伊人久久精品亚洲午夜| 亚洲第一区二区三区不卡| 久久99精品国语久久久| 日韩一本色道免费dvd| 久久久久精品久久久久真实原创| 大码成人一级视频| 亚洲成人手机| 51国产日韩欧美| 国产在线一区二区三区精| 成人午夜精彩视频在线观看| 国产精品.久久久| 少妇的逼水好多| 高清毛片免费看| 国产一区亚洲一区在线观看| 汤姆久久久久久久影院中文字幕| 插逼视频在线观看| 久久久久久久久久久免费av| 国产成人freesex在线| 午夜福利,免费看| 免费人成在线观看视频色| 久久精品国产亚洲av天美| 美女视频免费永久观看网站| 日韩亚洲欧美综合| 一区二区三区四区激情视频| 国产高清国产精品国产三级| 九色成人免费人妻av| 777米奇影视久久| 中国三级夫妇交换| 日韩中文字幕视频在线看片| 能在线免费看毛片的网站| 国产精品一二三区在线看| 特大巨黑吊av在线直播| 男人添女人高潮全过程视频| 一级毛片aaaaaa免费看小| 性高湖久久久久久久久免费观看| 精华霜和精华液先用哪个| 久久久久视频综合| 日韩欧美精品免费久久| 欧美日韩av久久| 亚洲不卡免费看| 久久久欧美国产精品| 99视频精品全部免费 在线| 街头女战士在线观看网站| 七月丁香在线播放| 国产男人的电影天堂91| 日本色播在线视频| 久久国产乱子免费精品| 精品久久久久久久久亚洲| 欧美少妇被猛烈插入视频| 一二三四中文在线观看免费高清| 亚洲成人一二三区av| 成年女人在线观看亚洲视频| 99re6热这里在线精品视频| 亚洲国产毛片av蜜桃av| 看十八女毛片水多多多| 肉色欧美久久久久久久蜜桃| 一个人看视频在线观看www免费| 天堂中文最新版在线下载| 久久国产乱子免费精品| www.av在线官网国产| 国产色婷婷99| 亚洲经典国产精华液单| 又黄又爽又刺激的免费视频.| 伦精品一区二区三区| a级毛片在线看网站| 三上悠亚av全集在线观看 | 国产一区二区在线观看日韩| 少妇裸体淫交视频免费看高清| 国产一区亚洲一区在线观看| 王馨瑶露胸无遮挡在线观看| 超碰97精品在线观看| 国产精品人妻久久久影院| 日韩强制内射视频| 一本久久精品| 免费观看无遮挡的男女| 黄色视频在线播放观看不卡| 久久这里有精品视频免费| 观看免费一级毛片| 蜜桃在线观看..| 亚洲激情五月婷婷啪啪| 国产真实伦视频高清在线观看| 国内揄拍国产精品人妻在线| 国国产精品蜜臀av免费| 免费久久久久久久精品成人欧美视频 | 久久国内精品自在自线图片| 一级,二级,三级黄色视频| 国产精品一区二区性色av| 国产熟女午夜一区二区三区 | 亚洲色图综合在线观看| 国产精品伦人一区二区| 黄色配什么色好看| 最黄视频免费看| 欧美日韩一区二区视频在线观看视频在线| 精品久久国产蜜桃| 国产精品国产三级国产专区5o| 久久精品国产鲁丝片午夜精品| 国产高清不卡午夜福利| 久久国产乱子免费精品| 国产精品免费大片| 国产成人aa在线观看| 久久精品国产a三级三级三级| 人人妻人人看人人澡| av播播在线观看一区| 2018国产大陆天天弄谢| 卡戴珊不雅视频在线播放| av在线app专区| 日日摸夜夜添夜夜添av毛片| 一级,二级,三级黄色视频| 久久久久人妻精品一区果冻| 春色校园在线视频观看| 欧美成人午夜免费资源| 最近中文字幕2019免费版| 黄片无遮挡物在线观看| 大片电影免费在线观看免费| 国产精品国产三级国产专区5o| 久久精品熟女亚洲av麻豆精品| 色5月婷婷丁香| 高清黄色对白视频在线免费看 | 欧美精品高潮呻吟av久久| 色5月婷婷丁香| 五月伊人婷婷丁香| 午夜免费鲁丝| 熟女人妻精品中文字幕| 国产在线视频一区二区| 国产极品天堂在线| 99热网站在线观看| 亚洲精品日本国产第一区| 99久久人妻综合| 精品视频人人做人人爽| 日本色播在线视频| 日韩强制内射视频| 中文字幕制服av| 亚洲欧美成人精品一区二区| 欧美区成人在线视频| 爱豆传媒免费全集在线观看| 91精品国产国语对白视频| 国产黄色视频一区二区在线观看| 久久毛片免费看一区二区三区| 美女脱内裤让男人舔精品视频| 国产真实伦视频高清在线观看| 少妇猛男粗大的猛烈进出视频| 99热全是精品| 黄色欧美视频在线观看| 在线观看免费日韩欧美大片 | 18禁裸乳无遮挡动漫免费视频| 热99国产精品久久久久久7| 亚洲美女搞黄在线观看| 国产亚洲av片在线观看秒播厂| 高清黄色对白视频在线免费看 | 少妇精品久久久久久久| 色5月婷婷丁香| 哪个播放器可以免费观看大片| 一边亲一边摸免费视频| 国产精品久久久久成人av| 亚洲第一av免费看| 国产在线免费精品| 亚洲美女黄色视频免费看| 一级毛片久久久久久久久女| 性色av一级| 嘟嘟电影网在线观看| 国产中年淑女户外野战色| 大片电影免费在线观看免费| av播播在线观看一区| 中文欧美无线码| 少妇的逼好多水| 青春草亚洲视频在线观看| 3wmmmm亚洲av在线观看| 日韩伦理黄色片| 最近最新中文字幕免费大全7| 最黄视频免费看| 一级毛片电影观看| 男男h啪啪无遮挡| 成人无遮挡网站| 精品一区在线观看国产| 久久久久久伊人网av| 97在线视频观看| 欧美bdsm另类| 日本与韩国留学比较| 蜜臀久久99精品久久宅男| 少妇人妻 视频| 国产精品三级大全| 免费久久久久久久精品成人欧美视频 | 日韩精品免费视频一区二区三区 | 久久精品国产鲁丝片午夜精品| 精品国产国语对白av| 久久久久久久精品精品| 777米奇影视久久| 午夜久久久在线观看| 大片电影免费在线观看免费| 久久人人爽人人爽人人片va| 欧美国产精品一级二级三级 | 欧美成人精品欧美一级黄| 欧美xxxx性猛交bbbb| 午夜激情福利司机影院| 乱人伦中国视频| 国产探花极品一区二区| 色网站视频免费| 精品熟女少妇av免费看| 边亲边吃奶的免费视频| 一区在线观看完整版| 街头女战士在线观看网站| 色吧在线观看| 国产 一区精品| 欧美日韩在线观看h| 日本av免费视频播放| 一本—道久久a久久精品蜜桃钙片| 亚洲自偷自拍三级| 国产一区亚洲一区在线观看| 六月丁香七月| 黄色视频在线播放观看不卡| 99热全是精品| 一级av片app| 18禁动态无遮挡网站| 大又大粗又爽又黄少妇毛片口| 中文资源天堂在线| 国产精品一区二区三区四区免费观看| 精品酒店卫生间| 99国产精品免费福利视频| 国产精品人妻久久久久久| 国产探花极品一区二区| 插逼视频在线观看| 美女大奶头黄色视频| 日韩人妻高清精品专区| 啦啦啦在线观看免费高清www| 国产黄片视频在线免费观看| 在现免费观看毛片| 日韩一区二区三区影片| 亚洲精品国产av蜜桃| 亚洲无线观看免费| 久久国产精品大桥未久av | 亚洲欧美一区二区三区黑人 | 午夜福利影视在线免费观看| 色视频www国产| 亚洲精品国产av蜜桃| 纯流量卡能插随身wifi吗| av又黄又爽大尺度在线免费看| 久久久久久久久大av| 又大又黄又爽视频免费| 国产熟女午夜一区二区三区 | 人妻一区二区av| 亚洲成色77777| 男人狂女人下面高潮的视频| 一本久久精品| 国产视频内射| 久久狼人影院| 大陆偷拍与自拍| 亚洲精品亚洲一区二区| 久久影院123| 9色porny在线观看| 亚洲精品色激情综合| 精品一区在线观看国产| 久久午夜福利片| 水蜜桃什么品种好| 又黄又爽又刺激的免费视频.| 国产精品不卡视频一区二区| 曰老女人黄片| 一本久久精品| a级毛色黄片| 国产深夜福利视频在线观看| 久久精品国产a三级三级三级| 丝袜喷水一区| 国产免费又黄又爽又色| 人人妻人人看人人澡| av有码第一页| 国产av码专区亚洲av| 在线观看av片永久免费下载| 99热全是精品| 久久精品久久精品一区二区三区| 国产真实伦视频高清在线观看| 日韩中文字幕视频在线看片| 99九九在线精品视频 | 老司机影院成人| 麻豆乱淫一区二区| 久久精品久久久久久久性| 乱码一卡2卡4卡精品| 国产成人精品无人区| 国产精品无大码| av福利片在线| 亚洲欧洲日产国产| 人人妻人人添人人爽欧美一区卜| 中文乱码字字幕精品一区二区三区| 精品一品国产午夜福利视频| 亚洲性久久影院| 国产永久视频网站| 亚州av有码| 久久精品久久精品一区二区三区| 成年人午夜在线观看视频| 国产美女午夜福利| 丝袜脚勾引网站| 亚洲不卡免费看| 三上悠亚av全集在线观看 | 人妻一区二区av| 国产有黄有色有爽视频| 免费大片黄手机在线观看| 交换朋友夫妻互换小说| 看十八女毛片水多多多| 亚洲国产日韩一区二区| 久久国产乱子免费精品| 日本爱情动作片www.在线观看| 欧美日韩av久久| 亚洲美女搞黄在线观看| 亚洲丝袜综合中文字幕| 免费看不卡的av| 国产黄片美女视频| 日韩av在线免费看完整版不卡| 精品国产露脸久久av麻豆| 久久人妻熟女aⅴ| 自拍偷自拍亚洲精品老妇| 国产黄频视频在线观看| 最黄视频免费看| 欧美国产精品一级二级三级 | 特大巨黑吊av在线直播| 精品一区二区三区视频在线| 国产69精品久久久久777片| 国产欧美亚洲国产| 天天操日日干夜夜撸| 99久久精品热视频| 黑人猛操日本美女一级片| 亚洲国产精品一区三区| 久久韩国三级中文字幕| 国产精品熟女久久久久浪| 亚洲精品国产av蜜桃| 18禁动态无遮挡网站| 在线观看免费视频网站a站| 一个人看视频在线观看www免费| 国产精品一区二区性色av| 国产欧美日韩综合在线一区二区 | 黄色欧美视频在线观看| 岛国毛片在线播放| 国产黄频视频在线观看| 十分钟在线观看高清视频www | 中文在线观看免费www的网站| 黑丝袜美女国产一区| 亚洲综合色惰| 国产精品熟女久久久久浪| 色5月婷婷丁香| 搡女人真爽免费视频火全软件| 丰满饥渴人妻一区二区三| 性高湖久久久久久久久免费观看| 一级片'在线观看视频| 久久久久久久久久久丰满| 在线观看av片永久免费下载| 免费观看性生交大片5| 熟女av电影| 精品国产乱码久久久久久小说| 人妻系列 视频| 中国美白少妇内射xxxbb| 91aial.com中文字幕在线观看| 男男h啪啪无遮挡| 中文字幕久久专区| 日韩中字成人| 在线免费观看不下载黄p国产| 国产一区二区三区综合在线观看 | 亚洲三级黄色毛片| 天堂中文最新版在线下载| 免费观看性生交大片5| 又爽又黄a免费视频| 最新中文字幕久久久久| 亚洲欧美日韩另类电影网站| 国产又色又爽无遮挡免| 夫妻午夜视频| 久久久久久久大尺度免费视频| 日日啪夜夜撸| 热99国产精品久久久久久7| 大又大粗又爽又黄少妇毛片口| 国产精品国产三级国产av玫瑰| 国产亚洲精品久久久com| 伦理电影免费视频| 国产永久视频网站| 成人国产av品久久久| 亚洲av不卡在线观看| 亚洲精品一区蜜桃| 少妇高潮的动态图| 一区在线观看完整版| 国产成人精品福利久久| 91久久精品电影网| 啦啦啦在线观看免费高清www| 欧美精品亚洲一区二区| 国产亚洲av片在线观看秒播厂| av专区在线播放| 大香蕉久久网| 桃花免费在线播放| 免费大片18禁| 国产黄色视频一区二区在线观看| 18禁动态无遮挡网站| 中国美白少妇内射xxxbb| 男人舔奶头视频| 日韩中文字幕视频在线看片| 少妇人妻精品综合一区二区| 十八禁高潮呻吟视频 | 久久国产乱子免费精品| 青春草亚洲视频在线观看| 一个人免费看片子| 熟女电影av网| 在线观看三级黄色| 老司机影院成人| 欧美激情极品国产一区二区三区 | 一边亲一边摸免费视频| 国产真实伦视频高清在线观看| 97超视频在线观看视频| 国产高清有码在线观看视频| 80岁老熟妇乱子伦牲交| 欧美精品一区二区大全| 18禁在线无遮挡免费观看视频| 久久精品国产a三级三级三级| 自拍偷自拍亚洲精品老妇| 一本久久精品| 欧美一级a爱片免费观看看| 亚洲美女视频黄频| 狂野欧美白嫩少妇大欣赏| 国产精品人妻久久久久久| 日韩三级伦理在线观看| 中文字幕精品免费在线观看视频 | 日韩av不卡免费在线播放| 丰满乱子伦码专区| 久久精品久久久久久久性| 边亲边吃奶的免费视频| 亚洲国产精品一区三区| 丰满饥渴人妻一区二区三| www.av在线官网国产| 欧美另类一区| 日韩一区二区三区影片| 伦精品一区二区三区| 看非洲黑人一级黄片| 久久久久国产精品人妻一区二区| 亚洲国产日韩一区二区| 日韩大片免费观看网站| 高清欧美精品videossex| 性高湖久久久久久久久免费观看| 日日撸夜夜添| 国产片特级美女逼逼视频| 国产爽快片一区二区三区| 日韩,欧美,国产一区二区三区| 看非洲黑人一级黄片| 亚洲成人手机| 欧美日韩av久久| 少妇高潮的动态图| 成人国产麻豆网| 欧美日韩国产mv在线观看视频| 一个人免费看片子| 亚洲欧美日韩东京热| 有码 亚洲区| 青春草国产在线视频| 成年人免费黄色播放视频 | 亚洲精华国产精华液的使用体验| 欧美精品人与动牲交sv欧美| 最后的刺客免费高清国语| 一个人免费看片子| 欧美精品人与动牲交sv欧美| 免费av中文字幕在线| 国产亚洲av片在线观看秒播厂| 国产乱来视频区| 亚洲av在线观看美女高潮| 日韩精品有码人妻一区| 国产日韩欧美亚洲二区| 亚洲欧洲国产日韩| 97精品久久久久久久久久精品| 最近中文字幕2019免费版| 嫩草影院新地址| 夫妻性生交免费视频一级片| 一二三四中文在线观看免费高清| 久久久久国产网址| 一区二区三区四区激情视频| 女性生殖器流出的白浆| 中文字幕av电影在线播放| 成人午夜精彩视频在线观看| 免费观看性生交大片5| 中文资源天堂在线| 亚洲色图综合在线观看| 99re6热这里在线精品视频| 97超视频在线观看视频| 又爽又黄a免费视频| 女人精品久久久久毛片| 乱人伦中国视频| 国产熟女欧美一区二区| 国产精品一区www在线观看| 欧美xxⅹ黑人| 少妇熟女欧美另类| 男人舔奶头视频| 婷婷色av中文字幕| 蜜臀久久99精品久久宅男| 亚洲高清免费不卡视频| 成年美女黄网站色视频大全免费 | 成人18禁高潮啪啪吃奶动态图 | 伦理电影大哥的女人| 99久久精品国产国产毛片| 人妻制服诱惑在线中文字幕| 热99国产精品久久久久久7| 九草在线视频观看| 五月天丁香电影| 大陆偷拍与自拍| 欧美精品亚洲一区二区| 性高湖久久久久久久久免费观看| 久久久久网色| a级毛片在线看网站| 特大巨黑吊av在线直播| av播播在线观看一区| 亚洲国产精品国产精品| 日本wwww免费看| 精品一品国产午夜福利视频| 在线观看三级黄色| 色5月婷婷丁香| 久久影院123| 波野结衣二区三区在线| 成人黄色视频免费在线看| 99热全是精品| 国产一区二区在线观看av| 人妻系列 视频| 极品人妻少妇av视频| 18禁动态无遮挡网站| 五月玫瑰六月丁香| 成人美女网站在线观看视频| 国产高清国产精品国产三级| 少妇被粗大的猛进出69影院 | 一级黄片播放器| 亚洲av日韩在线播放| 一级av片app| 亚洲精品日韩av片在线观看| 亚洲第一av免费看| 久久久国产一区二区| 国产探花极品一区二区| 最近中文字幕高清免费大全6| av.在线天堂| 国产亚洲精品久久久com| 蜜臀久久99精品久久宅男| 性色av一级| 成人毛片a级毛片在线播放| 亚洲国产欧美日韩在线播放 | 美女内射精品一级片tv| 性高湖久久久久久久久免费观看| 亚洲av免费高清在线观看| 天堂中文最新版在线下载| av国产精品久久久久影院| 欧美最新免费一区二区三区| 91精品国产国语对白视频| 永久免费av网站大全| 老司机影院成人| 观看av在线不卡| 纯流量卡能插随身wifi吗| 五月玫瑰六月丁香| 中文字幕久久专区| 日日爽夜夜爽网站| 精品酒店卫生间| 水蜜桃什么品种好| 亚洲精品亚洲一区二区| 十八禁高潮呻吟视频 | 国产欧美日韩精品一区二区| 激情五月婷婷亚洲| 国产高清不卡午夜福利| 99九九在线精品视频 | 欧美日韩国产mv在线观看视频| 一级,二级,三级黄色视频| 亚洲欧美一区二区三区国产| 最新的欧美精品一区二区| 99热这里只有是精品在线观看| 精品一区二区免费观看| 观看免费一级毛片| 日韩不卡一区二区三区视频在线| 蜜桃久久精品国产亚洲av| 国产精品成人在线| 久久国产精品男人的天堂亚洲 | 国产精品99久久久久久久久| 99九九线精品视频在线观看视频| 永久网站在线| 大香蕉97超碰在线| av福利片在线| 国产在线一区二区三区精| 久久久久视频综合| 成人综合一区亚洲| 2018国产大陆天天弄谢| 三上悠亚av全集在线观看 | 午夜精品国产一区二区电影| 国产极品天堂在线| 建设人人有责人人尽责人人享有的| 男女边摸边吃奶| 久久久久久久大尺度免费视频| 街头女战士在线观看网站| 如何舔出高潮| 黄片无遮挡物在线观看| 久久久久久久亚洲中文字幕| 大码成人一级视频| 国产精品一区二区三区四区免费观看| 久久国产亚洲av麻豆专区| 午夜日本视频在线| 国产精品久久久久久精品电影小说| 天天操日日干夜夜撸| 人妻系列 视频| 日本免费在线观看一区| 99久久精品国产国产毛片| 我的老师免费观看完整版| 我要看黄色一级片免费的| 22中文网久久字幕| 黑人巨大精品欧美一区二区蜜桃 | 人人妻人人爽人人添夜夜欢视频 | 日韩 亚洲 欧美在线| 国产精品秋霞免费鲁丝片| 内地一区二区视频在线| 精品一区在线观看国产| kizo精华| 夜夜骑夜夜射夜夜干| 国产精品久久久久久久电影| 亚洲国产欧美日韩在线播放 | 久久久a久久爽久久v久久| 99热网站在线观看| 亚洲欧美日韩东京热| 日日爽夜夜爽网站| 国产极品天堂在线| 午夜激情福利司机影院| 免费不卡的大黄色大毛片视频在线观看| 久久久a久久爽久久v久久| 国精品久久久久久国模美| 成人免费观看视频高清| 韩国高清视频一区二区三区|