亚洲免费av电影一区二区三区,日韩爱爱视频,51精品视频一区二区三区,91视频爱爱,日韩欧美在线播放视频,中文字幕少妇AV,亚洲电影中文字幕,久久久久亚洲av成人网址,久久综合视频网站,国产在线不卡免费播放

        ?

        基于GPU的高吞吐量QC-LDPC碼編碼器實現(xiàn)

        2021-03-12 06:12:56段運德
        關(guān)鍵詞:編碼方案碼率碼字

        段運德,黎 勇

        (1.重慶郵電大學(xué) 通信與信息工程學(xué)院,重慶 400065;2.重慶大學(xué) 計算機學(xué)院,重慶 400044)

        0 引 言

        1962年,R.Gallager[1]在他的博士論文中第一次提出了低密度奇偶校驗(low-density parity-check, LDPC)碼。當(dāng)時由于計算機技術(shù)并不發(fā)達(dá),LDPC碼譯碼復(fù)雜度高,導(dǎo)致其優(yōu)秀的性能未被發(fā)現(xiàn)。直到1996年D.J.C.Mackay等[2]重新研究LDPC碼,發(fā)現(xiàn)其性能可以接近香農(nóng)限,LDPC碼才重新開始受到重視,并迅速成為信道編碼研究中的重要方向。QC-LDPC碼是一類具有準(zhǔn)循環(huán)結(jié)構(gòu)的LDPC碼,因為其循環(huán)特性可以降低編譯碼復(fù)雜度,目前已被應(yīng)用在很多標(biāo)準(zhǔn)和系統(tǒng)中,比如IEEE 802.11n(WLAN),IEEE 802.16e(WiMAX)[3],量子秘鑰分發(fā)[4],ATSC3.0[5]等。2017年,3GPP標(biāo)準(zhǔn)會議將QC-LDPC碼確定為5G通信系統(tǒng)的長碼標(biāo)準(zhǔn)[6]。

        因為實際系統(tǒng)往往選擇QC-LDPC碼,而不是其他LDPC碼,因此,提高QC-LDPC碼的編碼速率有重要的實際意義。為提高編碼效率,大家普遍選擇現(xiàn)場可編程邏輯門陣列(field programmable gate array, FPGA)、專有電路(application-specific integrated circuit, ASIC)等硬件解決方案[7],硬件方案通常只針對特定碼型或碼率的LDPC碼有效,通用性差。相比而言,圖形處理單元(graphics processing unit, GPU)加速是一種更便宜、更靈活、更高效的軟件方案。本文通過把QC-LDPC碼的校驗矩陣轉(zhuǎn)換為準(zhǔn)循環(huán)的生成矩陣,再基于編碼過程中的并行部分和GPU的內(nèi)存、線程結(jié)構(gòu),給出一種吉比特速率的軟件編碼方案。該方案不僅通過單個碼字編碼并行化以及多個碼字同時編碼來提高編碼吞吐量,而且通過對片上常量內(nèi)存和共享內(nèi)存的合理使用進一步加快了編碼速率。仿真結(jié)果表明,相較于已有的GPU編碼方案,本文的方案通用性更強,編碼吞吐量比文獻(xiàn)[8]高;相較于其他硬件快速編碼方案,本文的編碼器在通用性和編碼吞吐量上均明顯高于文獻(xiàn)[9]的FPGA編碼器和文獻(xiàn)[10]的CMOS編碼器。

        1 通過校驗矩陣計算QC-LDPC碼準(zhǔn)循環(huán)結(jié)構(gòu)的生成矩陣

        QC-LDPC碼的校驗矩陣由具有準(zhǔn)循環(huán)特性的子矩陣組成,表示為

        (1)

        (1)式中:c

        rank(H)=rank(D)=r

        (2)

        設(shè)得到的D矩陣為

        (3)

        H中選擇q列后,剩下的t-q列構(gòu)成矩陣B,表示為

        (4)

        將H按列重排得到Hqc=[B|D],Hqc同樣為該QC-LDPC碼的校驗矩陣。設(shè)該碼的生成矩陣為Gqc,則有

        (5)

        (5)式中:Gqc為k×n階矩陣;n為該QC-LDPC碼碼長,n=tb;k為信息位長,k=tb-r;O是cb×(tb-r)階零矩陣。由(5)式和文獻(xiàn)[11]中的算法即可解出矩陣Gqc,且Gqc有如下結(jié)構(gòu)

        (6)

        (7)

        Q=

        (8)

        (7)—(8)式中:I為b階單位陣;O為b階零矩陣;Gi,j為b階方陣,且具有準(zhǔn)循環(huán)結(jié)構(gòu),即Gi,j中每行向量是其前一行向量向右循環(huán)移動一位得到;Oi,j為全零矩陣;qrow代表有qrow行子矩陣。Oi,j,Qi,j列數(shù)都為b,在計算Gqc時可以得到Oi,j和Qi,j的行數(shù)。本文先將Qi,j的行數(shù)保存在數(shù)組qNum[qrow]中以備后面程序使用,表示為

        qNum[qrow]={qNum0,qNum1,…,qNumqrow}

        (9)

        (9)式中,qNumi對應(yīng)Q的第i行子矩陣Qi,·的行數(shù)。這樣得到的Qi,j同樣具有準(zhǔn)循環(huán)結(jié)構(gòu),即Qi,j中每行向量是其前一行向量向右循環(huán)移動一位得到的。設(shè)gi,j為Gi,j的首行向量,qi,j為Qi,j的首行向量,則通過所有g(shù)i,j和qi,j的循環(huán)移位即可以得到矩陣G和Q,從而得到Gqc。這里,定義和gi,jqi,j為Gqc的生成向量。

        通過上述計算得到的生成矩陣Gqc通常具有(6)式的結(jié)構(gòu),即由子矩陣G和Q組成。但如果給定的QC-LDPC碼的校驗矩陣H是非奇異矩陣,且q=c,此時通過上述計算得到的Gqc只由子矩陣G組成,即Gqc=[G]。

        在求解Gqc的過程中,需要多次求GF(2)中矩陣的秩。為了加快求秩速度,本文采用了文獻(xiàn)[12]中的稀疏矩陣高斯消去算法求秩。

        通過以上處理,從QC-LDPC碼的校驗矩陣H得到了具有準(zhǔn)循環(huán)結(jié)構(gòu)的生成矩陣Gqc。根據(jù)其循環(huán)特性,即由首行向量gi,j和qi,j循環(huán)移位可得到Gqc。因此,在編碼程序中,僅需存儲生成向量gi,j和qi,j,節(jié)省了存儲空間。

        2 QC-LDPC碼的GPU編碼

        基于GPU的并行算法設(shè)計通常是一項具有挑戰(zhàn)性的工作。必須慎重考慮線程的數(shù)量、內(nèi)存的大小以及寄存器等資源的限制。同時為了獲得正確的結(jié)果,必須注意線程的同步。CUDA(compute unied device architecture)是NVIDIA的GPU提供并行計算的平臺和編程接口,本文將基于CUDA平臺編程仿真。在設(shè)計編碼時,本文將從以下一些方面進行優(yōu)化。表1給出了下文敘述中用到的一些參數(shù)。

        CUDA核函數(shù)初始化比CPU函數(shù)初始化需要更多時間,因此編碼時,核函數(shù)將同時輸入N=32×s(s為正整數(shù))個信息序列,然后在運行一次核函數(shù)時同時對N個碼字進行編碼,以加快編碼速度。CUDA執(zhí)行核函數(shù)的調(diào)度單位為線程束,每個線程束包含32個線程,因此,選擇N為32的倍數(shù)可以加快線程調(diào)度。

        表1 文中重要參數(shù)

        本文將定義結(jié)構(gòu)體Bitset保存長度為b個比特的向量,如gi,j和qi,j向量。本文中Hqc,H,G等矩陣的子矩陣階數(shù)為b,也用該結(jié)構(gòu)體來保存。該結(jié)構(gòu)體偽代碼如下。

        //一個int變量長32比特

        const int g_intLen←32;

        //g_num為保存長為b的向量需要int的個數(shù)

        const int g_num←(b-1)/g_intLen+1;

        struct Bitset{

        int val[g_num];

        }

        CUDA常量內(nèi)存是一種64 KByte的只讀緩存,它提供比全局內(nèi)存更大的訪問帶寬。編碼時生成矩陣Gqc是在第1節(jié)中已經(jīng)求解出來的常量,不會改變,因此,將Gqc的生成向量gi,j和gi,j按圖1的格式存儲在常量內(nèi)存的Bitset類型的二維數(shù)組generatorMtx中。

        圖1中,t,q,qrow參數(shù)含義見表1。使用CUDA常量內(nèi)存存儲Gqc的二維數(shù)組代碼為

        __constant__ __device__

        Bitset generatorMtx[t-q+qrow][q];

        圖1 矩陣Gqc在常量內(nèi)存中的映射Fig.1 Map of Gqc stored in constant memory

        編碼算法的CUDA核函數(shù)原型為

        __global__

        voidcuEncode(int *mess, int *code);

        其中,mess是一個數(shù)組指針,保存N個信息序列,該數(shù)組包含N×k個比特。code為編碼輸出的碼字?jǐn)?shù)組指針,數(shù)組包含N×n個比特。

        當(dāng)調(diào)用CUDA核函數(shù)編碼時需要進行線程分配,本文將分配二維線程塊numThread,其維度為(t-q+qrow)×q,numThread和存儲Gqc的二維數(shù)組generatorMtx具有相同的維度。在同一個線程塊中,每個線程將同時讀取gi,j和qi,j,同時進行計算。為了充分利用GPU線程,本文還將定義一維的線程網(wǎng)格numBlock,網(wǎng)格包含N個線程塊,每個線程塊中進行單個碼字的編碼,核函數(shù)運行一次將同時編得N個碼字。調(diào)用該核函數(shù)的代碼為

        dim3 numThread(t-q+qrow, q), numBlock(N);

        cuEncode << > > (mess, code);

        此外,在核函數(shù)中,本文使用共享內(nèi)存用于一個線程塊中所有線程的通信,共享內(nèi)存的變量作為中間變量,存儲累加和,即將同一個線程塊中所有線程的計算結(jié)果累加到該變量中。由于核函數(shù)訪問共享內(nèi)存的速度比全局內(nèi)存快10倍,因此,使用共享內(nèi)存可以減少一個線程塊中所有線程的內(nèi)存訪問時間,加快編碼。

        以上是本文設(shè)計GPU編碼時,從變量存儲方面做的一些優(yōu)化。接下來,本文將實現(xiàn)GPU編碼核函數(shù)。設(shè)m為一個信息序列向量,相應(yīng)的碼字為c,由于QC-LDPC碼是一種線性分組碼,則有

        c=mGqc

        (10)

        根據(jù)前面的論述及(10)式,GPU編碼核函數(shù)的偽代碼如下。

        算法:GPU并行編碼。

        輸入:函數(shù)傳入N個信息序列保存在CUDA全局內(nèi)存數(shù)組mess[N×k]中。

        輸出:最后得到的N個碼字輸出到數(shù)組code[N×n]中。

        1)定義共享內(nèi)存數(shù)組,即__shared__ Bitset codTmp[q],并將其初始化為全零。

        2) __syncthreads(),同步所有線程塊。

        3)bTmp←generatorMtx[threadIdx.x][threadIdx.y]

        4)if threadIdx.x>= t-q then

        //計算第threadIdx.x個線程塊處理的信息

        //序列對應(yīng)的起始位置bitLoc,其中qNum[i]為

        //(9)式中得到的數(shù)據(jù),k為信息位長。

        for j←bitLoc to bitLoc + qNum[i]-1 do

        loc ← j + k × blockIdx.x;

        if mess[loc] == 1 then

        //異或相加

        codTmp[threadIdx.y]←

        codTmp[threadIdx.y]⊕bTmp;

        end if

        bTmp向右循環(huán)移動一位;

        end for

        else

        bitLoc ← threadIdx.x × b;

        //計算存放碼字的起始位置messLoc,

        //用于存放信息比特。如果校驗矩陣H滿秩,

        //則得到系統(tǒng)碼,否則碼字中僅前一部分

        //比特為信息比特,n為碼長。

        messLoc←threadIdx.x×b+n×blockIdx.x;

        for j←bitLoc to bitLoc + b-1 do

        loc←j+k×blockIdx.x;

        if mess[loc] == 1 then

        codTmp[threadIdx.y]←

        codTmp[threadIdx.y]⊕bTmp

        end if

        bTmp向右循環(huán)移動一位;

        if threadIdx.y == 0 then

        code[messLoc]←mess[loc];

        ++messLoc;

        end if

        end for

        end if

        5) __syncthreads(),同步所有線程塊。

        6)

        if threadIdx.x == 0 then

        messLoc←blockIdx.x×n+(t-q)×b

        +threadIdx.y×b;

        for i←b-1 to 0 do

        code[messLoc]←codTmp[threadIdx.y][i];

        ++messLoc;

        end for

        end if

        上述核函數(shù)中,參數(shù)的含義如表1。threadIdx,blockIdx為CUDA中索引當(dāng)前線程和當(dāng)前線程塊的變量。__syncthreads()為CUDA中同步所有線程的函數(shù)。

        調(diào)用上述核函數(shù)產(chǎn)生N個碼字保存在數(shù)組code[N·n],code數(shù)組中碼字的排列方式如圖2。

        圖2 數(shù)組中N個碼字的排列方式Fig.2 Arrange of the N codes in the array

        3 GPU編碼的仿真分析

        本文仿真采用Nvidia TITAN Xp顯卡平臺,使用英偉達(dá)官方工具集CUDA Toolkit 9.2,在Visual Studio 2017下編譯程序,平臺其他參數(shù)如表2。

        表2 仿真電腦配置

        實驗中測試了碼長從幾百到一萬多比特的幾種高碼率QC-LDPC碼。相應(yīng)校驗矩陣大小及Tanner圖的邊數(shù)如表3。

        表3 仿真用的QC-LDPC碼

        為了評估編碼算法的性能,本文計算了編碼吞吐量Kraw,其粗略計算式為

        (11)

        (11)式中:tim為對一個碼字編碼平均所用時間,單位為s;messLen為該QC-LDPC碼的信息位長,單位為bit。

        在進行GPU編碼時,本文為核函數(shù)分配N個線程塊,N為32的倍數(shù),運行一次核函數(shù)編得N個碼字。本文測試了隨著N的增加,GPU編碼吞吐量的變化情況,仿真結(jié)果如圖3。

        圖3 核函數(shù)分配不同線程塊后的編碼性能Fig.3 Encoding performance with kernel function assigneddifferent number of thread blocks

        從仿真結(jié)果來看,隨著N的增大,這幾種QC-LDPC碼的吞吐量在前期增加較快,但在N大約增加到20×32后,編碼吞吐量不再增加,且在一個范圍中波動。這是由于前期GPU存在很多空閑資源,隨著N的增大,分配給核函數(shù)的線程增多,核函數(shù)運行一次同時編得的碼字增加,吞吐量增加,編碼效率明顯提高。但后期,GPU中使用的常量內(nèi)存、全局內(nèi)存、共享內(nèi)存等資源達(dá)到上限,因此,編碼吞吐量停止增加。

        實驗中也將傳統(tǒng)的LDPC碼近似下三角矩陣編碼算法[13]與該GPU加速的編碼算法進行了對比。該傳統(tǒng)算法是一種串行的CPU編碼算法,算法復(fù)雜度為O(n),實驗結(jié)果如表4。

        表4 與傳統(tǒng)的編碼算法對比

        從表4可以看出,本文提出的GPU編碼算法對于3種不同長度的QC-LDPC碼,其編碼吞吐量均超過了10 Gbit/s,傳統(tǒng)編碼算法最高吞吐量為12.1 Mbit/s,遠(yuǎn)不及本文GPU加速的算法。

        目前GPU加速的LDPC碼譯碼器很多,但編碼器不多?,F(xiàn)有的LDPC碼快速編碼方案大多是基于FPGA、集成電路的硬件編碼方案,接下來本文將分別討論。

        4 同其他GPU加速的編碼器對比

        在文獻(xiàn)[14]中,作者提出了針對DVB-S2系統(tǒng)中的LDPC碼的GPU編碼器。該LDPC碼的校驗矩陣具有圖4中的形態(tài)結(jié)構(gòu),其中,黑點處為非零元素,空白處為零元素。具有這種校驗矩陣結(jié)構(gòu)的DVB-S2 LDPC碼在設(shè)計GPU編碼器時可以節(jié)省更多的存儲空間,因而可以同時對更多碼字進行編碼。

        在同樣的Nvidia TITAN Xp顯卡平臺上,文獻(xiàn)[14]中的編碼器對碼長16 200,碼率8/9的DVB-S2 LDPC碼一次性可以編8 000個碼字,編碼吞吐量可以達(dá)到20 Gbit/s;由圖3可知,本文提出的GPU編碼器同時對480個(12 769, 12 320)QC-LDPC碼字進行編碼時,吞吐量為10.174 Gbit/s;最多可同時編碼1 344個碼字,此時的吞吐量約為8.0 Gbit/s。雖然本文提出的GPU編碼器編碼速度沒有文獻(xiàn)[14]快,但本文針對的是一般QC-LDPC碼的編碼方案,而文獻(xiàn)[14]針對的是DVB-S2 LDPC碼這種特定碼型LDPC碼的編碼方案。

        文獻(xiàn)[8]提出了針對任意LDPC碼的GPU編碼方案。它假設(shè)LDPC碼生成矩陣為q×s維矩陣,然后直接簡單地分配q個GPU線程進行并行編碼。該方案沒有充分考慮對單個碼字編碼時的并行處理,導(dǎo)致其編碼速度不快。該文獻(xiàn)仿真GPU為Tegra K1,在和傳統(tǒng)的編碼算法[13]進行仿真對比時,文中給出了不同數(shù)量碼字同時編碼的運行時間曲線圖,它的編碼器僅比傳統(tǒng)的編碼算法快2~4倍。但從表4可以看出,本文提出的編碼器比傳統(tǒng)算法[13]快900~40 180倍。

        5 同硬件快速編碼器比較

        在同其他硬件快速編碼器對比時,本文選擇了2種較快的硬件編碼器:①基于CMOS技術(shù)的編碼器[10];②基于FPGA的編碼器[9]。

        文獻(xiàn)[10]中提出了一種基于130 nmCMOS技術(shù)實現(xiàn)的編碼器,是目前針對802.11ac標(biāo)準(zhǔn)的最快的一種CMOS編碼器。編碼時,該編碼器主要通過在QC-LDPC碼校驗矩陣的列方向上并行化和在電路中引入循環(huán)移位寄存器來提高編碼吞吐量。該文的仿真結(jié)果表明,對于標(biāo)準(zhǔn)中的(1 944, 1 620)QC-LDPC碼在100 MHz的頻率時可以達(dá)到最高7.7 Gbit/s的編碼吞吐量。

        為了便于對比,本文同樣選擇了802.11ac標(biāo)準(zhǔn)中的碼字進行仿真。802.11ac標(biāo)準(zhǔn)中規(guī)定了一系列具有特殊結(jié)構(gòu)的不同碼長、不同碼率的QC-LDPC碼。本文仿真選擇了該標(biāo)準(zhǔn)中1/2碼率的(648, 324)QC-LDPC碼和(1 944, 972)QC-LDPC碼,以及3/4碼率的(1 944, 1 458)QC-LDPC碼和5/6碼率的(1 944, 1 620)QC-LDPC碼,仿真結(jié)果如表5。

        表5 與CMOS編碼器對比

        文獻(xiàn)[10]的CMOS編碼器可以對802.11ac標(biāo)準(zhǔn)中的所有QC-LDPC碼進行編碼,但不能對其他QC-LDPC碼編碼。文獻(xiàn)[10]對標(biāo)準(zhǔn)中不同碼率的碼進行了仿真,但只給出了編碼最快的(1 944, 1 620)QC-LDPC碼的吞吐量數(shù)據(jù)。本文中的GPU編碼器可以對任意QC-LDPC碼編碼,本實驗中選擇了802.11ac中3種碼率的QC-LDPC碼進行仿真。從表5可以看出,本文的編碼器吞吐量最低為5 Gbit/s,最高為9.6 Gbit/s。且對于(1 944, 1 620)QC-LDPC碼,本文的編碼器比文獻(xiàn)[10]的CMOS編碼器提高了1.9 Gbit/s的吞吐量,效果明顯。

        FPGA是編碼領(lǐng)域中常用的硬件實現(xiàn)技術(shù),目前有很多LDPC碼的FPGA編譯碼器方案,但其中速率較快的方案同樣只針對特定的碼型或碼率的LDPC碼,不具有通用性。本文討論一種2017年提出的FPGA編碼方案[9],該方案可以對WIMAX標(biāo)準(zhǔn)(即IEEE 802.16標(biāo)準(zhǔn))中不同碼長、不同碼率的QC-LDPC碼進行快速編碼。

        本文選擇了WIMAX標(biāo)準(zhǔn)中1/2碼率的(1 536, 768)QC-LDPC碼和(2 304, 1 152)QC-LDPC碼,以及3/4碼率的(1 536, 1 152)QC-LDPC碼和(2 304, 1 728)QC-LDPC碼,文獻(xiàn)[9]仿真平臺為117 MHz的FPGA平臺,仿真結(jié)果如表6。

        表6 與FPGA編碼器對比

        文獻(xiàn)[9]的FPGA編碼器是目前針對WIMAX標(biāo)準(zhǔn)中所有QC-LDPC碼通用的最快的FPGA編碼器之一,但從表6的仿真結(jié)果來看,本文提出的GPU編碼器吞吐量是該FPGA編碼器的3.94~7.73倍。而且本文提出的編碼方案可以對任意QC-LDPC碼進行快速編碼,通用性更好。

        6 結(jié) 論

        本文首先引入QC-LDPC碼具有準(zhǔn)循環(huán)特性的生成矩陣,用這樣的生成矩陣編碼,其編碼過程中許多步驟可以并行化。同時,GPU具有高速并行運算的特點。因此,本文基于GPU加速,給出了一種針對QC-LDPC碼通用的并行編碼方案。編碼時,GPU一個線程塊中多個線程并行編碼得到一個碼字,本文方案分配了多個線程塊同時進行多個碼字的編碼。此外,在編碼過程中使用了共享內(nèi)存、常量內(nèi)存等高速片上存儲器,縮短了數(shù)據(jù)訪問時間,提高了編碼吞吐量。從仿真結(jié)果可以看出,目前無論是基于CMOS,F(xiàn)PGA等硬件技術(shù)的QC-LDPC碼編碼方案,還是已有的基于GPU或CPU的軟件編碼方案都有一定的局限性,它們局限于特殊的碼型、碼率等,通用性不好。本文提出的GPU編碼方案不僅可以針對不同碼型、不同碼率的QC-LDPC碼實現(xiàn)快速編碼,通用性好,而且編碼速度也有明顯的提升。因此,本文提出的GPU加速編碼器在實際的高速編碼系統(tǒng)中有很好的應(yīng)用前景。

        猜你喜歡
        編碼方案碼率碼字
        基于功能類別和技術(shù)參數(shù)的刀具編碼方案設(shè)計
        基于唯一標(biāo)識的ATP車載設(shè)備編碼方案研究
        放 下
        揚子江詩刊(2018年1期)2018-11-13 12:23:04
        基于改進粒子群算法的毫米波大規(guī)模MIMO混合預(yù)編碼方案
        數(shù)據(jù)鏈系統(tǒng)中軟擴頻碼的優(yōu)選及應(yīng)用
        基于狀態(tài)機的視頻碼率自適應(yīng)算法
        放下
        揚子江(2018年1期)2018-01-26 02:04:06
        基于場景突變的碼率控制算法
        X264多線程下碼率控制算法的優(yōu)化
        計算機工程(2015年8期)2015-07-03 12:19:56
        三種預(yù)編碼方案對OFDM系統(tǒng)峰均比的影響分析
        中國新通信(2015年9期)2015-05-30 16:17:07
        久久免费精品视频老逼| 欧洲日本一线二线三线区本庄铃| 国产亚洲欧美成人久久片| 视频精品熟女一区二区三区| 99久久精品一区二区国产| 精品国产免费一区二区三区 | 国产偷国产偷亚洲清高| 尤物蜜芽福利国产污在线观看 | 黄瓜视频在线观看| 久久精品国产99精品国偷| 日韩少妇高潮在线视频| 日韩三级一区二区不卡| 国产精品欧美一区二区三区不卡| 日韩AV无码免费二三区| 一区二区日本影院在线观看| 亚洲中文字幕av天堂自拍| 国产在线精品一区二区在线看| 日韩AV无码一区二区三| av免费网站不卡观看| 极品少妇xxxx精品少妇偷拍| 老熟妻内射精品一区| 免费在线观看一区二区| 国产高清人肉av在线一区二区| 欧美黑寡妇特a级做爰| 二区三区视频| 日韩精品高清不卡一区二区三区| 亚洲av无码国产精品色| 午夜福利92国语| 激情 一区二区| 午夜男女靠比视频免费| 久久久亚洲欧洲日产国码αv| 久久综合五月天| 中文字幕色一区二区三区页不卡 | 中国丰满熟妇xxxx性| 中文 国产 无码免费| 神马不卡影院在线播放| 人人妻人人澡人人爽国产一区| 在线观看免费a∨网站| 一区二区久久精品66国产精品| 亚洲av无码国产精品色| 久久久久亚洲av无码专区导航|