999精品在线视频,手机成人午夜在线视频,久久不卡国产精品无码,中日无码在线观看,成人av手机在线观看,日韩精品亚洲一区中文字幕,亚洲av无码人妻,四虎国产在线观看 ?

基于GPU的密碼S盒代數性質評估方法

2022-09-25 08:42:32蔡婧雯韋永壯劉爭紅
計算機應用 2022年9期

蔡婧雯,韋永壯*,劉爭紅

(1.廣西密碼學與信息安全重點實驗室(桂林電子科技大學),廣西桂林 541004;2.廣西無線寬帶通信與信號處理重點實驗室(桂林電子科技大學),廣西桂林 541004)

0 引言

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

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

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

1 預備知識

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

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

則差分均勻度定義為:

當差分均勻度越小,S 盒的差分分布更均勻,安全性更好。

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

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

當透明階越小時,抵抗差分功耗攻擊的能力越強,安全性越好。

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

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

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

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

2 CUDA的并行計算

2.1 GPU與CUDA

隨著人工智能、大數據等計算領域的不斷發展,計算復雜度越來越大,圖像處理器(GPU)從幫助CPU 做圖像和圖形運算轉至海量數據處理上,涉及云計算、生物計算、天文學等多個領域,逐步成為了計算領域的研究熱點。

由于CPU 與GPU 所應用的場景不同,兩者的架構大不相同。從圖1 中CPU 與GPU 的架構相比可知,CPU 與GPU 有如下區別:

圖1 CPU與GPU的架構區別Fig.1 Difference between CPU and GPU architectures

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

2)CPU 具有強大的控制邏輯單元,在運算過程中提供邏輯預測能力降低延時;而GPU 控制能力稍遜色于CPU,運算過程中可以將多個訪問合并成較少的訪問。

3)CPU 有大量的緩存空間降低計算延時,而GPU 只有少量的緩存空間,與CPU 的緩存空間功能不同,GPU 的緩存空間將線程所需要訪問的相同數據合并訪問動態隨機存取存儲器。

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

圖2 CPU與GPU任務耗時比較Fig.2 Task time consumption comparison of CPU and GPU

GPU 的造價和功耗與相同計算能力的CPU 相比,GPU 的造價和功耗相對較低。在計算領域中構建CPU 集群的超級計算機,造價昂貴。根據摩爾定律可以得知當CPU 計算速度達到一定程度時提升空間受限,GPU 的出現滿足了需要計算大量數據而無法使用巨型計算機的用戶需求。

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

2.2 線程塊的分配

為了發揮GPU 的最大并行計算效率,在執行內核函數過程中,需要合理配置線程塊數量及每塊線程塊中的線程數量。線程塊的數目由配置參數所劃分的網格所決定,通常為32 的倍數最佳。通過測試可知,將線程塊的線程數量為512時具有最大的計算能力。

2.3 并行計算的影響因素

總結前文所述的GPU 與CUDA 的計算特點,影響并行計算效率的因素可以總結為以下3 點:

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

2)減少訪問GPU 全局內存。為了減少內存訪問產生的時延和消耗,在GPU 中應盡量減少過多的跳躍式訪問,最大限度減少因對GPU 內存訪問而造成的延遲。

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

3 NBC算法

NBC 算法為中國密碼學會舉辦的全國密碼算法設計競賽分組算法[11]第二輪入選算法之一,其采用廣義Feistel 結構[12],算法加密共有三種模式,具體如表1 所示。

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

本文使用的算法是數據分組長度為128 比特、密鑰長度為128 比特的NBC 算法。設第i輪的輸入為Xi=輸出為NBC-128/128 算法結構如圖3 所示。

圖3 八分支的1輪NBC-128/128結構Fig.3 One-round NBC-128/128 structure with 8 branches

NBC-128/128 算法的S 盒采用16 級NFSR來構造16 比特S 盒,S 盒構造圖如圖4 所示。設S 盒的16 比特的輸入為S0S1…S15,當全體內部狀態經過迭代20 輪后形成S 盒輸出。

圖4 NBC-128算法的S盒構造Fig.4 S-box structure of NBC-128 algorithm

算法設計者稱構造出來的S 盒的差分均勻度Diff(S)=22,非線性度NL(S)=31 982,透明階RTO=15.982 6。

4 基于GPU的16比特密碼S盒代數性質評估

由于在CPU 下求解差分均勻度、非線性度及透明階的算法效率較低,在本章中,將傳統求解密碼S 盒代數性質評估方法進行優化,分別討論基于單GPU 模式和多GPU 模式下將內核函數切片至多線程中,實現多線程并行化計算。

4.1 單GPU對16比特密碼S盒性質評估

根據共享式內存的結構特點和對S 盒性質評估的求解流程,本文提出了單塊GPU 環境下的CPU-GPU 異構模式,并行架構如圖5 所示。

圖5 CPU-GPU異構并行流程Fig.5 CPU-GPU heterogeneous parallel flowchart

程序在運行時控制一塊GPU,創建多個線程共同完成計算任務。具體步驟如下所示:

1)檢測顯卡設備。函數cudaSetDevice()表示檢測主機設備的顯卡個數,當檢測到主機存在可使用的顯卡時,將對算法進行CUDA 并行計算做好準備;

2)讀取數據并復制入GPU。采用cudaMalloc()函數在設備Device 中開辟計算中所需要參數的空間。由于GPU 在計算過程中,無法直接讀取CPU 內存中的數據,故在計算前需要在設備Decive 開辟相應的空間。

3)當Device 中開辟了相應的空間大小后,使用cudaMemcpy()函數將所需要的參數S 盒復制進入GPU 內。

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

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

6)檢查并返回結果。當每一個線程完成了內核函數中的計算任務時,使用cudaGetLastError()函數檢查內核函數在計算過程中是否存在錯誤:若存在錯誤,將錯誤返回至CPU中;若不存在,利用函數cudaMemcpy()將計算結果返回至CPU 中,計算結束。

求解復雜度分析如下:

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

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

算法1 測試主程序。

輸入 S 盒;

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

4.2 多GPU對大狀態S盒性質評估

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

在使用多GPU 構架中,選擇單個節點連接到高速串行計算機擴展總線標準(Peripheral Component Interconnect express,PCIe)總線上,具體架構如圖6 所示。程序在運行時使用函數cudaSetDecice()對GPU 設備組上的各設備進行綁定,使得每個線程管理一個GPU,實現多個GPU 并行工作。與單GPU 結構相比,多GPU 結構可以開辟更多的線程,運算速度得到進一步提升。

圖6 多GPU節點架構Fig.6 Multi-GPU node architecture

由于同步傳輸的并行化計算中,傳輸數據占用了大量的時間。本節利用多流技術與異步技術相結合,在計算過程中使計算過程與數據傳輸兩個步驟進行重疊,從而減少一部分時間的開銷。有無重疊優化的時間開銷對比如圖7 所示。在無重疊優化時,由于默認只有一個流隊列,此時所有的計算過程皆為串行執行。先對數據傳輸至GPU 的全局內存內,傳輸完畢后再進行數據計算,等待GPU 內的所有線程計算完畢后再將結果復制回CPU 內。作為對比,在使用重疊技術進行時間優化后,當一個流隊列在計算部分數據的同時,另一個流隊列可以對剩下數據進行傳輸至GPU 內等待計算。當一個流上的數據計算完畢后,利用另一個流隊列傳輸回CPU 內,下一個流隊列等待數據傳輸。計算與傳輸時間重疊技術的優化既能保持計算任務仍按照串行執行,又能掩蓋GPU 與CPU 數據傳輸之間所帶來的大量時間開銷,從而進一步減少程序所需要的執行時間,提高并行效率。

圖7 有無重疊優化的時間開銷對比Fig.7 Comparison of time cost with and without overlapping optimization

利用多GPU 對大狀態S 盒進行評估過程具體如下:

1)在CPU 端獲取已有的GPU 設備數量和每個GPU 設備信息。利用CUDA 中自帶的函數cudaGetDeviceCount(&ngpus),讀取已有的GPU 設備數量,并將GPU 數量信息存儲在變量ngpus中,可通過設備號dev進行選擇使用GPU設備。

2)在同一節點上的GPU 設備構成GPU 設備組,GPU 設備組內的GPU 設備直接進行通信和數據傳輸。

3)在CPU 端進一步準備計算所需要的數據集,根據GPU設備組的數量,將數據平分至各GPU 設備上,另外在CPU 端設置S 盒,在CPU 端將S 盒以結構體的形式傳輸至GPU 設備組中對應的常量存儲區中,S 盒參數都在對應的GPU 設備運行過程中將會被核函數多次調用。

4)在CPU 端設置循環遍歷所有的GPU 設備,將GPU 設備組中的GPU 分別置于對應并行流上,通過對工作流在不同時間下的操作和阻塞實現GPU 設備的異步,設置CUDA 工作流Steam 的異步操作隱藏了部分訪問延遲和實現了任務的并發執行,減少了數據處理時間。

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

6)為了確保核函數運行時有較好的性能,延用上一節對GPU 的線程數分配,使用了128 個線程組,其中每個線程組512 個線程,共計65 536 個線程數。

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

8)當線程中循環遍歷完所有的塊,完成內核函數Kernel的計算后,利用重疊流的思想保證每個線程計算后優先傳輸至CPU 內。

9)當所有線程都計算完畢后,CPU 端對GPU 設備組的各GPU 設備返回的結果進行統一歸總,并按照規定的格式進行輸出。

算法2 多GPU 測試主程序。

輸入 S 盒;

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

5 測試與結果分析

5.1 測試環境

本文實驗環境所使用的CPU 為Intel Xeon Silver 4210 2.20 GHz;GPU 為NVIDIA Quadro RTX 8000;在多GPU 環境下,共使用4 塊相同型號的GPU,且顯卡型號為NVIDIA Quadro RTX 8000;操作系統為Ubuntu 18.04.4 LTS,64 bits;編程環境為CUDA 7.0、GCC 7.5.0。本實驗的CPU 代碼用的C 語言進行編寫,GPU 代碼用CUDA C 進行編寫。

5.2 測試結果

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

圖8 對比CPU、單塊GPU及多塊GPU下的運行時間Fig.8 Comparison of running time under CPU,single GPU and multi-GPU

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

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

6 結語

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

主站蜘蛛池模板: 毛片视频网址| 91色爱欧美精品www| 欧美天天干| 国产精品天干天干在线观看| a级毛片毛片免费观看久潮| 国产69精品久久久久孕妇大杂乱| 日韩在线1| 国产综合网站| 亚洲精品国产乱码不卡| 成人久久精品一区二区三区| 久久成人免费| 日韩a级片视频| 在线精品亚洲一区二区古装| 色男人的天堂久久综合| 在线精品亚洲一区二区古装| 亚洲精品中文字幕无乱码| 国产精品成人啪精品视频| 亚洲中文字幕97久久精品少妇| 国产白浆在线| 亚洲日韩高清在线亚洲专区| 国产丰满成熟女性性满足视频 | 精品91视频| 特级毛片免费视频| 欧美精品成人一区二区视频一| 99热这里只有精品免费| 成人国产小视频| 国产在线自在拍91精品黑人| 午夜少妇精品视频小电影| 免费毛片在线| 国产福利一区视频| 国产91成人| 国产精品lululu在线观看| 国产va欧美va在线观看| 99免费视频观看| 午夜三级在线| 国产第一页第二页| 亚卅精品无码久久毛片乌克兰| 国产精品区视频中文字幕| 亚洲热线99精品视频| 婷婷午夜影院| 欧美福利在线播放| hezyo加勒比一区二区三区| 2021最新国产精品网站| 91日本在线观看亚洲精品| 亚洲日韩精品欧美中文字幕| 91色国产在线| 亚洲人成网站在线播放2019| 国产一国产一有一级毛片视频| 日韩精品资源| 亚洲第一视频区| 中文字幕乱码中文乱码51精品| 亚洲av日韩av制服丝袜| 亚洲无码熟妇人妻AV在线| 在线欧美日韩| 香港一级毛片免费看| 中文字幕无码制服中字| 国产第二十一页| 日韩视频福利| 成人日韩视频| 国内嫩模私拍精品视频| 久草国产在线观看| 中国一级特黄大片在线观看| 全裸无码专区| 欧美激情首页| 99在线视频免费| 尤物精品国产福利网站| 中文无码精品a∨在线观看| 日韩在线视频网| 亚洲欧美日韩综合二区三区| 国产美女无遮挡免费视频| 亚洲国产成人精品无码区性色| 国产精品女熟高潮视频| 国产在线91在线电影| 国产高清免费午夜在线视频| 91视频99| 强乱中文字幕在线播放不卡| 国产成人无码AV在线播放动漫| 欧美成a人片在线观看| 无码一区18禁| 亚洲欧美不卡中文字幕| 粉嫩国产白浆在线观看| 国产jizzjizz视频|