色哟哟视频在线观看-色哟哟视频在线-色哟哟欧美15最新在线-色哟哟免费在线观看-国产l精品国产亚洲区在线观看-国产l精品国产亚洲区久久

0
  • 聊天消息
  • 系統(tǒng)消息
  • 評論與回復(fù)
登錄后你可以
  • 下載海量資料
  • 學(xué)習(xí)在線課程
  • 觀看技術(shù)視頻
  • 寫文章/發(fā)帖/加入社區(qū)
會員中心
电子发烧友
开通电子发烧友VIP会员 尊享10大特权
海量资料免费下载
精品直播免费看
优质内容免费畅学
课程9折专享价
創(chuàng)作中心

完善資料讓更多小伙伴認(rèn)識你,還能領(lǐng)取20積分哦,立即完善>

3天內(nèi)不再提示

gemv優(yōu)化總結(jié)

perfxlab ? 來源:澎峰科技PerfXLab ? 2023-05-25 09:08 ? 次閱讀
加入交流群
微信小助手二維碼

掃碼添加小助手

加入工程師交流群

有朋友來信說:

1. “除了以NVIDIA(英偉達(dá))為例,能不能談點(diǎn)國產(chǎn)GPU優(yōu)化的經(jīng)驗(yàn)分享?”

2. “老講國外的東西,你們能不能支持一下國產(chǎn)CPU和加速卡?“

這里解釋一下原因:

1. N卡的資料和環(huán)境大家都比較好找,對于學(xué)習(xí)GPU并行優(yōu)化編程的朋友比較友善。

2. 暫時受限于商業(yè)保密,我們相信后續(xù)會逐步開放起來,學(xué)習(xí)的平臺和環(huán)境也容易找到。到時就可以分享一些國產(chǎn)CPU和加速卡的優(yōu)化經(jīng)驗(yàn)出來。

------ 正文分割線 ------

本文主要是介紹如何對gemv算法進(jìn)行優(yōu)化。gemv,即矩陣向量乘,即計算一個矩陣A與一個向量x的乘積,這是并行計算中的經(jīng)典話題。個人感覺,gemv的優(yōu)化核心是需要考慮不同shape的情況,然后針對型地進(jìn)行優(yōu)化。本篇文章會先介紹一下針對不同shape設(shè)計不同的并行算法,然后說明一下優(yōu)化思路和相關(guān)優(yōu)化技巧,最后說一下實(shí)驗(yàn)效果,在A矩陣列數(shù)為16 128的時候,我寫的gemv能擁有超越cublas的性能表現(xiàn)。

一、前言

首先介紹一下gemv算法。給定矩陣A和向量x,gemv需要計算兩者的乘積,示意圖如下:

6a2366c2-fa8c-11ed-90ce-dac502259ad0.png

gemv

二、針對不同shape的并行算法設(shè)計

這次講到并行算法設(shè)計,什么叫并行算法設(shè)計。每個人的理解都不太一樣,在GPU中,我的理解就是:設(shè)計block和thread的workload,說白了就是要搞清楚一個block負(fù)責(zé)哪部分的計算,一個thread要負(fù)責(zé)哪部分的計算。而設(shè)計的原則就是盡可能地減少訪存,提高數(shù)據(jù)的復(fù)用概率,然后讓所有的處理器都滿負(fù)荷地進(jìn)行工作,不能浪費(fèi)。

2.1 針對n=32

對于n=32的情況,我們將每個block設(shè)置為256個線程,4個warp,然后每個warp負(fù)責(zé)一行元素的計算。每個warp要對x進(jìn)行訪問,然后在warp內(nèi)部進(jìn)行一次reduce求和操作。

6a38baf4-fa8c-11ed-90ce-dac502259ad0.jpg

n=32

代碼如下:

template 
__device__ __forceinline__ float warpReduceSum(float sum) {
    if (WarpSize >= 32)sum += __shfl_down_sync(0xffffffff, sum, 16); // 0-16, 1-17, 2-18, etc.
    if (WarpSize >= 16)sum += __shfl_down_sync(0xffffffff, sum, 8);// 0-8, 1-9, 2-10, etc.
    if (WarpSize >= 8)sum += __shfl_down_sync(0xffffffff, sum, 4);// 0-4, 1-5, 2-6, etc.
    if (WarpSize >= 4)sum += __shfl_down_sync(0xffffffff, sum, 2);// 0-2, 1-3, 4-6, 5-7, etc.
    if (WarpSize >= 2)sum += __shfl_down_sync(0xffffffff, sum, 1);// 0-1, 2-3, 4-5, etc.
    return sum;
}

// if N == 32
__global__ void Sgemv_v0( 
    float * __restrict__ A,
    float * __restrict__ x,
    float * __restrict__ y, 
    const int M,
    const int N) {
    // Block index
    int bx = blockIdx.x;

    // Thread index
    int tx = threadIdx.x;
    int ty = threadIdx.y;

    const int warp_size=32;
    int laneId= tx % warp_size;
    int current_row = blockDim.y * bx + ty;

    if(current_row < M){
        float res=0;
        int kIteration = N/warp_size;
        if(kIteration==0) kIteration=1;
        #pragma unroll
        for(int i=0; i< kIteration; i++){
            int current_col = i*warp_size + laneId;
            res += A[current_row*N + current_col] * x[current_col];
        }
        res = warpReduceSum(res);
        if(laneId==0) y[current_row]=res;
    }
}

2.2 針對n=128

對于n=128的情況,同樣讓warp負(fù)責(zé)一行元素的計算,但是因?yàn)槊啃械脑乇容^多,所以采用了float4進(jìn)行向量化的訪存。能夠有更高的訪存效率。

6a53326c-fa8c-11ed-90ce-dac502259ad0.jpg

n=128

代碼如下:

template 
__device__ __forceinline__ float warpReduceSum(float sum) {
    if (WarpSize >= 32)sum += __shfl_down_sync(0xffffffff, sum, 16); // 0-16, 1-17, 2-18, etc.
    if (WarpSize >= 16)sum += __shfl_down_sync(0xffffffff, sum, 8);// 0-8, 1-9, 2-10, etc.
    if (WarpSize >= 8)sum += __shfl_down_sync(0xffffffff, sum, 4);// 0-4, 1-5, 2-6, etc.
    if (WarpSize >= 4)sum += __shfl_down_sync(0xffffffff, sum, 2);// 0-2, 1-3, 4-6, 5-7, etc.
    if (WarpSize >= 2)sum += __shfl_down_sync(0xffffffff, sum, 1);// 0-1, 2-3, 4-5, etc.
    return sum;
}

// if N>= 128
__global__ void Sgemv_v1( 
    float * __restrict__ A,
    float * __restrict__ x,
    float * __restrict__ y, 
    const int M,
    const int N) {
    // Block index
    int bx = blockIdx.x;

    // Thread index
    int tx = threadIdx.x;
    int ty = threadIdx.y;

    const int warp_size=32;
    int laneId= tx % warp_size;
    int current_row = blockDim.y * bx + ty;

    if(current_row < M){
        float res=0;
        int kIteration = (N/warp_size)/4;
        if(kIteration==0) kIteration=1;
        A = &A[current_row*N];
        #pragma unroll
        for(int i=0; i< kIteration; i++){
            int current_col_vec = (i*warp_size + laneId);
            float4 current_val= reinterpret_cast(A)[current_col_vec];
            float4 current_x = reinterpret_cast(x)[current_col_vec];
            res += current_val.x*current_x.x;
            res += current_val.y*current_x.y;
            res += current_val.z*current_x.z;
            res += current_val.w*current_x.w;
        }
        res = warpReduceSum(res);
        if(laneId==0) y[current_row]=res;
    }
}

2.3 針對n=16

對于n=16的情況,讓一個warp負(fù)責(zé)兩行元素的計算。以warp0為例,0-15號線程負(fù)責(zé)第0行元素的計算,而16-31號線程負(fù)責(zé)第1行元素的計算。

6a6a3796-fa8c-11ed-90ce-dac502259ad0.jpg

n=16

代碼如下:

template 
__device__ __forceinline__ float warpReduceSum(float sum) {
    if (WarpSize >= 32)sum += __shfl_down_sync(0xffffffff, sum, 16); // 0-16, 1-17, 2-18, etc.
    if (WarpSize >= 16)sum += __shfl_down_sync(0xffffffff, sum, 8);// 0-8, 1-9, 2-10, etc.
    if (WarpSize >= 8)sum += __shfl_down_sync(0xffffffff, sum, 4);// 0-4, 1-5, 2-6, etc.
    if (WarpSize >= 4)sum += __shfl_down_sync(0xffffffff, sum, 2);// 0-2, 1-3, 4-6, 5-7, etc.
    if (WarpSize >= 2)sum += __shfl_down_sync(0xffffffff, sum, 1);// 0-1, 2-3, 4-5, etc.
    return sum;
}

// if N <= 16
template <
    const int ROW_PER_WARP
    > 
__global__ void Sgemv_v2( 
    float * __restrict__ A,
    float * __restrict__ x,
    float * __restrict__ y, 
    const int M,
    const int N) {
    // Block index
    int bx = blockIdx.x;

    // Thread index
    int tx = threadIdx.x;
    int ty = threadIdx.y;

    const int warp_size=32;
    int laneId= tx % warp_size;
    int current_warp_row = (blockDim.y * bx + ty) * ROW_PER_WARP;
    const int kWarp_size = warp_size / ROW_PER_WARP;
    int kLaneId = laneId % kWarp_size;
    int current_thread_row = current_warp_row + laneId / kWarp_size;

    if(current_thread_row < M){
        float res=0;
        int current_col = kLaneId;
        res += A[current_thread_row * N + current_col] * x[current_col];
        res = warpReduceSum(res);
        if(kLaneId==0) y[current_thread_row]=res;
    }
}

三、優(yōu)化思路:

上一節(jié)說明了如何針對不同維度的n進(jìn)行優(yōu)化,這一節(jié)說明一下為什么要這么設(shè)計,以及這樣的設(shè)計方式能夠帶來什么樣的好處。主要考慮的因素有兩個,如下:

3.1 盡可能地讓warp中的32個線程忙碌

這個主要是針對n<32的情況,例如n=16,如果使用一個warp來負(fù)責(zé)一行元素的計算,那么warp中有一半的元素都是浪費(fèi)的。所以讓一個warp來負(fù)責(zé)多行元素的計算,這樣讓32個線程全部忙碌起來。

3.2 盡可能地提高訪存效率

① global mem->register

將數(shù)據(jù)從global memory搬運(yùn)到寄存器上時,最重要的就是考慮是不是進(jìn)行了合并訪存。在這里,我們只考慮矩陣數(shù)據(jù)在global mem中是地址對齊的,即n是2的多次冪。上述的三種并行實(shí)現(xiàn)中,warp中的32個線程都是連續(xù)地訪問32個float或者128個float,因而滿足了合并訪存的條件,確保了global -> register的訪存效率。

② shared mem->register

說到這里,可能會有讀者好奇,上述的代碼都沒有用到shared mem。為啥要說這個點(diǎn)。我們可以再仔細(xì)看看上述的三種并行實(shí)現(xiàn),以第2種為例,一個block中有4個warp,每個warp都需要對x進(jìn)行一次global上的訪存,所以一個block有4次訪存。如果將x存儲到shared mem中,4個warp都去訪問shared mem上的x,這樣的話,對于global的訪存就從4次變成1次。直觀上會有性能提升,但不幸的是,如果用shared mem的話,將global mem的數(shù)據(jù)搬運(yùn)至shared mem需要有同步操作,這又會導(dǎo)致性能的下降。總的來說,使用shared mem并沒有得到顯著的提升,不過還是在這里說明一下。

③ 向量化訪存

向量化訪存就是一個老生常談的話題了,說白了就是盡可能地使用128bit的訪存指令,這個在reduce、sgemm、elementwise專題上說了很多,就不再多說。

四、實(shí)驗(yàn)與總結(jié)

筆者在V100上進(jìn)行了實(shí)驗(yàn),迭代1000次,用nsight進(jìn)行了測試,性能數(shù)據(jù)如下:

sgemv M N my_sgemv time(ns) cublas(ns) my_sgemv/cublas
v0 16384 32 10341 8386 81.1%
v1 16384 128 14284 15848 110.9%
v2 16384 16 6903 7576 109.7%

可以看出,在n=16以及n=128的情況下,都比cublas性能要好。n=32的情況要差于cublas。如果再加上向量化訪存應(yīng)該能夠有更好的性能表現(xiàn)。由于我實(shí)在沒時間再進(jìn)行深入,有心的同學(xué)可以改改代碼看看效果 :)。

審核編輯:彭靜
聲明:本文內(nèi)容及配圖由入駐作者撰寫或者入駐合作網(wǎng)站授權(quán)轉(zhuǎn)載。文章觀點(diǎn)僅代表作者本人,不代表電子發(fā)燒友網(wǎng)立場。文章及其配圖僅供工程師學(xué)習(xí)之用,如有內(nèi)容侵權(quán)或者其他違規(guī)問題,請聯(lián)系本站處理。 舉報投訴
  • cpu
    cpu
    +關(guān)注

    關(guān)注

    68

    文章

    11113

    瀏覽量

    218204
  • 存儲
    +關(guān)注

    關(guān)注

    13

    文章

    4553

    瀏覽量

    87711
  • 編程
    +關(guān)注

    關(guān)注

    88

    文章

    3692

    瀏覽量

    95498
  • 澎峰科技
    +關(guān)注

    關(guān)注

    0

    文章

    74

    瀏覽量

    3417

原文標(biāo)題:深入淺出GPU優(yōu)化系列:gemv優(yōu)化

文章出處:【微信號:perfxlab,微信公眾號:perfxlab】歡迎添加關(guān)注!文章轉(zhuǎn)載請注明出處。

收藏 0人收藏
加入交流群
微信小助手二維碼

掃碼添加小助手

加入工程師交流群

    評論

    相關(guān)推薦
    熱點(diǎn)推薦

    HBase性能優(yōu)化方法總結(jié)

    hbase響應(yīng)速度;9. 避免出現(xiàn)region熱點(diǎn)現(xiàn)象,啟動按照table級別進(jìn)行balance。以上是對HBase性能優(yōu)化方法的概要總結(jié),有HBase性能優(yōu)化需求的,可以在此基礎(chǔ)上延伸學(xué)習(xí),會有一定收獲的!
    發(fā)表于 04-20 17:16

    30種SQL語句優(yōu)化總結(jié)

    必須掌握的30種SQL語句優(yōu)化
    發(fā)表于 04-21 11:38

    改善深層神經(jīng)網(wǎng)絡(luò)--超參數(shù)優(yōu)化、batch正則化和程序框架 學(xué)習(xí)總結(jié)

    《深度學(xué)習(xí)工程師-吳恩達(dá)》02改善深層神經(jīng)網(wǎng)絡(luò)--超參數(shù)優(yōu)化、batch正則化和程序框架 學(xué)習(xí)總結(jié)
    發(fā)表于 06-16 14:52

    文檔基于DSP的視頻監(jiān)控系統(tǒng)的優(yōu)化仿真的設(shè)計總結(jié)文檔

    該文檔為基于DSP的視頻監(jiān)控系統(tǒng)的優(yōu)化仿真的設(shè)計總結(jié)文檔,是一份很不錯的參考資料,具有較高參考價值,感興趣的可以下載看看………………針對目前IGBT驅(qū)動電路復(fù)雜的缺點(diǎn),本文以德國西門康公司
    發(fā)表于 08-27 16:19

    嵌入式Java虛擬機(jī)優(yōu)化技術(shù)總結(jié)的太棒了

    嵌入式Java虛擬機(jī)優(yōu)化技術(shù)總結(jié)的太棒了
    發(fā)表于 04-25 06:47

    總結(jié)下電機(jī)控制中對程序算法優(yōu)化的辦法

    (用到了三角函數(shù))都比較消耗電機(jī)主控芯片的計算能力。在考慮算法實(shí)現(xiàn)的時候,都需要針對主控芯片的實(shí)際性能進(jìn)行一定優(yōu)化,才能確保算法能夠順利運(yùn)行。這里我總結(jié)下電機(jī)控制中對程序算法優(yōu)化的辦法。數(shù)據(jù)的概念浮點(diǎn)數(shù)
    發(fā)表于 08-27 06:37

    電機(jī)控制中對程序算法優(yōu)化的辦法總結(jié)

    (用到了三角函數(shù))都比較消耗電機(jī)主控芯片的計算能力。在考慮算法實(shí)現(xiàn)的時候,都需要針對主控芯片的實(shí)際性能進(jìn)行一定優(yōu)化,才能確保算法能夠順利運(yùn)行。這里我總結(jié)下電機(jī)控制中對程序算法優(yōu)化的辦法。數(shù)據(jù)的概念...
    發(fā)表于 09-07 06:19

    GPRS優(yōu)化思路總結(jié)報告

    GPRS優(yōu)化思路總結(jié)報告:一、概述 2二、無線優(yōu)化的思路 2三、(E)GPRS網(wǎng)絡(luò)資源容量分析優(yōu)化 53.1、(E)GPRS網(wǎng)絡(luò)拓?fù)浣Y(jié)構(gòu) 63.
    發(fā)表于 07-27 21:29 ?26次下載

    DSP程序優(yōu)化總結(jié)

    DSP程序優(yōu)化總結(jié)
    發(fā)表于 10-23 14:24 ?2次下載
    DSP程序<b class='flag-5'>優(yōu)化</b><b class='flag-5'>總結(jié)</b>

    區(qū)塊鏈共識算法的效能優(yōu)化研究及總結(jié)

    。由于共識算法的資源花銷、能源耗費(fèi)以及性能之間相互關(guān)聯(lián)且關(guān)系復(fù)雜,因此有必要從¨效能”的角度對現(xiàn)有區(qū)塊鏈的共識算法加以分析,并總結(jié)研究思路。文中總結(jié)了區(qū)塊鏈共識算法的效能優(yōu)化研究進(jìn)展。首先定義區(qū)塊鏈共識算法的效
    發(fā)表于 04-25 11:35 ?4次下載
    區(qū)塊鏈共識算法的效能<b class='flag-5'>優(yōu)化</b>研究及<b class='flag-5'>總結(jié)</b>

    TD-LTE網(wǎng)絡(luò)優(yōu)化經(jīng)驗(yàn)總結(jié)解析

    TD-LTE網(wǎng)絡(luò)優(yōu)化經(jīng)驗(yàn)總結(jié)解析說明。
    發(fā)表于 04-27 10:30 ?23次下載

    DC-DC電源系統(tǒng)的優(yōu)化設(shè)計總結(jié)

    DC-DC電源系統(tǒng)的優(yōu)化設(shè)計總結(jié)(電源技術(shù)期刊咋樣)-該文檔為DC-DC電源系統(tǒng)的優(yōu)化設(shè)計總結(jié)文檔,是一份不錯的參考資料,感興趣的可以下載看看,,,,,,,,,,,,,,,,,
    發(fā)表于 09-22 11:45 ?26次下載
    DC-DC電源系統(tǒng)的<b class='flag-5'>優(yōu)化</b>設(shè)計<b class='flag-5'>總結(jié)</b>

    深入淺出GPU優(yōu)化系列:gemv優(yōu)化

    這次講到并行算法設(shè)計,什么叫并行算法設(shè)計。每個人的理解都不太一樣,在GPU中,我的理解就是:設(shè)計block和thread的workload,說白了就是要搞清楚一個block負(fù)責(zé)哪部分的計算,一個thread要負(fù)責(zé)哪部分的計算。
    的頭像 發(fā)表于 05-25 09:03 ?3189次閱讀
    深入淺出GPU<b class='flag-5'>優(yōu)化</b>系列:<b class='flag-5'>gemv</b><b class='flag-5'>優(yōu)化</b>

    總結(jié)FasterTransformer Encoder優(yōu)化技巧

    FasterTransformer BERT 包含優(yōu)化的 BERT 模型、高效的 FasterTransformer 和 INT8 量化推理。
    的頭像 發(fā)表于 05-30 15:15 ?1674次閱讀
    <b class='flag-5'>總結(jié)</b>FasterTransformer Encoder<b class='flag-5'>優(yōu)化</b>技巧

    性能優(yōu)化之路總結(jié)

    針對老項(xiàng)目,去年做了許多降本增效的事情,其中發(fā)現(xiàn)最多的就是接口耗時過長的問題,就集中搞了一次接口性能優(yōu)化。本文將給小伙伴們分享一下接口優(yōu)化的通用方案。 ? ? 一、接口優(yōu)化方案總結(jié) 1
    的頭像 發(fā)表于 06-17 15:00 ?581次閱讀
    主站蜘蛛池模板: 国产亚洲papapa | 色网站在线免费观看 | 天干天干天干夜夜爽av | 极品尤物在线观看 | 久久精品国产成人 | 精品国精品国产自在久国产应用 | 国产正在播放 | 中文字幕在线国产 | 97久久天天综合色天天综合色hd | 成人爽a毛片在线视频 | 97久久精品人人澡人人爽 | 少妇精品 | 欧美人与动物xxxxx | 92av视频| cao在线 | 亚洲春色av无码专区在线播放 | 日本网站在线看 | 午夜1000集 | 特级毛片全部免费播放器 | 夜夜操网 | 亚洲精品视频一二三区 | 久久99精品久久久久久水蜜桃 | 性视频欧美 | 青青操在线观看视频 | 色综合婷婷 | 国产传媒中文字幕 | 色老板精品凹凸在线视频观看 | 亚洲tv久久爽久久爽 | 亚洲最大中文字幕 | 天堂色播| 国产精品一区二区av | 一个人看的免费高清www视频 | 成人看片17c.com | 精品蜜臀av在线天堂 | 久久天堂综合亚洲伊人hd妓女 | 国产精品久久久久久亚洲影视公司 | 国产小视频自拍 | 无码精品人妻一区二区三区人妻斩 | 性折磨bdsm欧美激情另类 | 肉嫁高柳在线 | 亚洲午夜福利在线视频 | 日产一二三四五六七区麻豆 | 阿v天堂在线观看 | 亚洲国产综合无码一区 | 一区二区三区日韩视频 | 欧美成人午夜 | 黄a大片 | 少妇翘臀亚洲精品av图片 | 国内女人喷潮完整视频 | 最新中文字幕在线观看视频 | 日本老妇70sex另类 | 日日碰狠狠躁久久躁9 | 日韩女同疯狂作爱系列5 | 中出白浆| 欧美黄色免费在线观看 | 欧美人伦禁忌dvd放荡欲情 | 在线不卡二区 | 99久久无码一区人妻a黑 | 亚洲成人免费看 | 亚洲中文字幕aⅴ天堂 | 日本一本久 | 激情九九 | 免费毛片小视频 | 中文字幕亚洲一区二区va在线 | 91精品国模一区二区三区 | av在线大全 | 日韩在线精品 | 5151精品国产人成在线观看 | 国产精品成人免费一区久久羞羞 | 国产欧美久久一区二区三区 | 精品女同一区二区三区 | 美女日日日 | 久久久一本精品99久久精品66 | 粉嫩av一区二区在线播 | 久久久久女教师免费一区 | 13女裸体慰在线观看 | 久久精品99国产精品酒店日本 | 亚洲日韩国产av无码无码精品 | 国产视频一二区 | 一区二区三区人妻无码 | 日本黄色成人 | va婷婷在线免费观看 | 亚洲色图欧美日韩 | 亚洲中文字幕无码中文字在线 | 国产在线观看a | 亚洲第一av网站 | 一区二区日韩欧美 | 久久99精品久久久久久久久久久久 | 狂野欧美性猛交xxxxx视频 | 欧美一区二区三区在线看 | 国产亚洲欧美日韩俺去了 | 亚洲国产aⅴ成人精品无吗 欧洲熟妇色xxxx欧美老妇多毛网站 | 亚洲欧洲一区二区在线观看 | 色吊丝永久性观看网站免费 | 天天干天天操心 | 蜜桃视频一区二区在线观看 | 久久精品视频亚洲 | 91精品国产乱码久久久久 | 六月丁香婷婷综合 | 亚洲一区二区在线播放 | 性欧美videos另类艳妇3d | 男人的天堂色 | 国产成人在线一区 | 中国a级大片 | 亚洲r成人av久久人人爽澳门赌 | jizz欧美性11| 欧美精品乱人伦久久久久久 | 日韩精品在线播放 | 激性欧美激情在线 | 婷婷综合激情 | 窝窝午夜看片 | 九色视频导航 | 火车卧铺高h肉辣文虐 | 私人成片免费观看 | 91不卡视频| 狠狠色综合网站久久久久久久高清 | av福利站| 中国女人内谢69xxxx | 欧美日本免费一区二区三区 | 97超碰免费在线 | 免费观看成人鲁鲁鲁鲁鲁视频 | 欧美性aaa | 久久久精品视频网站 | 亚洲天堂不卡 | 亚洲一区二区三区视频在线 | 亚洲成人免费视频 | 最近日韩免费视频 | 成人精品网站在线观看 | 性做久久久久 | 久久国产精品视频一区 | 国偷自产视频一区二区久 | 一区二区三区国产在线 | 成人看片网 | 91亚洲国产成人久久精品麻豆 | 日本理论中文字幕 | 亚洲第一色图 | 91精品国产综合久久精品 | 午夜精品久久久久久久四虎 | 亚洲欧洲日本无在线码 | 得得啪在线| 免费激情网址 | 成人免费在线网站 | 国产成人歌舞艳r舞 | 怡红院免费的全部视频 | 一区二区在线播放视频 | 夜夜撸av| 日韩av在线网| 中文无码精品一区二区三区 | 四虎永久在线精品免费观看网站 | 国产精品欧美亚洲777777 | 美女一区二区三区网av | 久久久久久久久久久影院 | www.狠狠操.com | 国产在线播放一区 | 激情视频一区二区三区 | 久久久精品99 | 精产国品一区二区三区四区 | 色爽交| 最近中文字幕mv在线资源 | 狠狠躁天天躁综合网 | 久久久久久久久久久免费精品 | 日韩在线www | 户外少妇对白啪啪野战 | 狠狠躁天天躁中文字幕 | 欧美视频黄色 | 中文字幕在线观看免费 | 亚洲精品久久久中文字幕 | 欧洲成人在线 | 国产欧美日韩精品a在线观看 | 中文字幕av日韩精品一区二区 | 亚洲熟妇国产熟妇肥婆 | 97性无码区免费 | 国产性xxxxx| 摸丰满大乳奶水www免费 | 天天干天天看 | 嫩草社区 | 国内一级黄色 | 欧美人妻精品一区二区三区 | 在线播放五十路熟妇 | 大桥未久亚洲无av码在线 | 日韩深夜影院 | 999视频在线播放 | 热99re久久精品这里都是精品 | 精品国产aⅴ无码一区二区 亚洲人成人无码网www国产 | 亚洲videos| 风间由美一区 | 伊人蕉 | 国产又白又嫩又紧又爽18p | 日本丰满大乳mm | 欧美成人一区二区三区在线观看 | 亚洲不卡av不卡一区二区 | 久久人妻内射无码一区三区 | а√天堂资源官网在线资源 | 午夜精品久久久久久99热 | 白嫩丰满少妇xxxxx性张津瑜 | 欧美性生活精品 | 91丨九色丨蝌蚪丨老版 | 99国产精品自在自在久久 | www蜜臀| 久久久久久久久免费看无码 | 91精品专区 | 78亚洲精品久久久蜜桃网 | 少妇一级淫免费放 | 精品国产一区二区三区久久久狼 | 三级自拍| 日韩人妻无码一区二区三区综合部 | 久草视频手机在线观看 | 国产农村妇女aaaaa视频 | 78国产伦精品一区二区三区 | 丁香桃色午夜亚洲一区二区三区 | 欧美日韩一区二区视频在线观看 | 久久久精品成人免费观看 | 中文字幕亚洲精品日韩一区 | 农村人伦偷精品视频a人人澡 | 天天夜夜骑 | 亚洲天堂网在线观看 | 亚洲免费看片 | 亚洲精品综合五月久久小说 | 婷婷五综合 | 久久精品网站视频 | 欧美最新精品videossexohd | 日韩在线视频网站 | 婷婷色九月 | 欧美一区中文字幕 | 国产高清精品软件 | 国产精品久久久久久久久久软件 | 欧美成人片在线观看 | 免费一级片网址 | 糖心av | 美女黄色真播 | 人人舔人人干 | 淫欲的代价k8经典网 | 午夜av剧场 | 深夜福利麻豆 | 日韩视频成人 | 破了亲妺妺的处免费视频国产 | 粉嫩欧美一区二区三区高清影视 | 亚洲女人av久久天堂 | 国产下药迷倒白嫩丰满美女j8 | 国产无遮挡又黄又爽在线观看 | 国产大片黄 | 热播之家 | 免费看av大片 | 1024中文字幕| 在线看国产 | 毛片a久久99亚洲欧美毛片 | 全黄一级片 | 欧洲精品视频在线 | 自拍偷拍第 | 无码专区亚洲综合另类 | 国产做受入口竹菊 | 国产伦理片在线观看 | 91精品国产乱码久久久久 | 欲妇荡岳丰满少妇岳 | 黄色天堂网 | 国产精品久久无码一区二区三区网 | 意大利少妇愉情理伦片 | 亚洲一区二区三区四区五区不卡 | 国产精品视频yy9299一区 | 五月天婷婷缴情五月免费观看 | 午夜国产一区二区 | 国产一精品一av一免费 | 精品国产乱码一区二区三 | 中文有码一区 | 妺妺窝人体色www聚色窝 | 亚洲一区二区三区写真 | 玖玖在线| 亚洲午夜久久久影院 | 999精品在线视频 | 欧美另类极品videosbes | 日本www视频 | 大尺度做爰床戏呻吟色戒韩国 | 黄色成人免费观看 | 久久久久久av无码免费网站 | 精品国产乱码久久久久夜深人妻 | 国产精品1000 | 一本大道东京热无码 | 欧美视频中文在线看 | 夜夜操夜夜 | jizjiz中国少妇高潮水多 | 亚洲淫片| 国产亚洲福利 | 丁香花五月 | 伊人情人色综合网站 | 一本加道在线 | 最近免费中文字幕中文高清百度 | 天天射日日操 | 午夜美女在线 | 亚洲性生活视频 | 91小视频 | 日本精品久久久久中文字幕 | 亚洲另类春色 | 伊人久久狼人 | 亚洲你懂得 | 国产免费网 | 亚洲专区第一页 | 欧美精品高清 | 激情内射亚洲一区二区三区爱妻 | 在线视频se | 18禁裸乳无遮挡啪啪无码免费 | 91小宝寻花一区二区三区 | 亚洲福利精品 | 中文字幕高清在线 | 国产乱码精品一区二区三区中文 | 免费日韩av| 果冻传媒mv免费播放在线观看 | 欧美深度肠交惨叫 | 欧美不卡一二三 | 美女胸18大禁视频网站 | 成人免费国产 | 国产无区一区二区三麻豆 | aa性欧美老妇人牲交免费 | 欧美在线视频免费播放 | h在线网站 | 新版本天堂资源在线中文8的特点 | 日本人视频69式jzzij | 日韩欧美成人精品 | 四虎4hu永久免费深夜福利 | 女人内谢aaaa免费视频 | 在线看片wwwzzz | 视频黄色免费 | xx中文字幕乱偷avxx | 国产主播av在线 | 亚洲va欧美va天堂v国产桃 | 男人的亚洲天堂 | 农村妇女做爰偷拍视频 | 最新日韩中文字幕 | 无码日韩精品一区二区人妻 | 午夜精品久久久久久久99芒果 | 天堂视频在线 | 波多一区二区 | 夜色jjj.av| 久草网在线观看 | 亚洲一区精品在线 | 日韩大片免费观看视频播放 | 无码无套少妇毛多69xxx | 亚洲91在线| 最新av不卡 | 国产精欧美一区二区三区 | 欧美日韩一区二区三区在线观看视频 | 国产精品一二区 | 午夜福利理论片在线观看 | 国语自产拍精品香蕉在线播放 | 好吊色视频988gao在线观看 | 亚洲国产精品激情在线观看 | 国产精品久久久久久久久久新婚 | 天天射一射 | 亚洲 另类 在线 欧美 制服 | 精品久久久久久一区二区 | 狠狠躁天天躁夜夜添人人 | 国产精品99久久久久久久vr | 日韩av一区在线观看 | h成人在线观看 | 日韩成人在线网站 | 国产日本免费 | 白嫩丰满少妇xxxxx性视频 | 免费观看性欧美大片无片 | 欧美一区日韩一区 | 国产三级久久久精品麻豆三级 | 大肉大捧一进一出视频出来呀 | 日韩国产精品免费 | 亚洲色婷婷综合开心网 | 久久国产精品99国产精 | 午夜精品久久久久久久久久久久久 | 亚洲午夜久久久精品一区二区三区 | 1769国产精品 | 欧美丰满熟妇xx猛交 | 超碰青草 | 韩国免费a级毛片 | 日本精品视频一区二区三区 | 一区二区国产精品精华液 | 青青操在线 | 欧美视频在线一区二区三区 | 国产成人精品日本亚洲专区61 | 国产久操视频 | 狠狠色综合欧美激情 | 日本私人影院 | 人人澡人人妻人人爽人人蜜桃麻豆 | 男女啪啪做爰高潮免费网站 | 乱码av| 欧美性大战久久久久xxx | 欧美成人一区二区三区片免费 | 在线中文字幕播放 | 国产亚洲va天堂va777 | 狠狠干欧美 | 免费99精品国产自在在线 | 久草中文在线 | 免费观看久久 | 色婷婷av一区二区三区软件 | 在线欧美激情 | 丁香婷婷成人 | 超碰97人人草| 国产精品白浆无码流出 | 国产香蕉在线观看 | 久久综合av色老头免费观看 | 国产精品久久伊人 | 大学生久久香蕉国产线看观看 | 亚洲日韩欧美视频 | 韩国日本在线 | 欧美大片在线看免费观看 | 亚洲综合色区中文字幕 | 亚洲免费精品网站 | 日本韩国三级 | 丰满少妇xbxb毛片日本视频 | aⅴ在线视频男人的天堂 | 黑人操亚洲美女 | 九色视频国产 | 18禁裸男晨勃露j毛免费观看 | 国产一级美女视频 | 108种啪姿势大全动态图 | 人妻饥渴偷公乱中文字幕 | 亚洲福利在线观看 | yy6080理aa级伦大片一级 | 一级大片儿 | 久久蜜桃资源一区二区老牛 | 古风h啪肉禁欲 | 人妻人人澡人人添人人爽人人玩 | 国产一级特黄aaa大片评分 | 国产成人高潮免费观看精品 | 精品久久久久久中文字幕人妻最新 | 北京少妇宾馆露脸对白 | 久久婷婷五月综合色国产香蕉 | 国语对白精品 | 黑桃tv视频一区二区 | 特级全黄久久久久久久久 | 波多野结衣视频播放 | 99精品成人| 久久12| 精品亚洲一区二区三区四区五区 | 欧洲性网站 | 外国一级片 | 亚洲欧洲视频在线 | 日韩国产精品视频 | 黄色av免费 | 亚洲精品www久久久 亚洲精品www久久久久久 | 丰满岳乱妇一区二区 | 成年人免费看的视频 | 国产精品亚州 | 九九热在线精品视频 | 国产视频你懂得 | 久久亚洲私人国产精品 | 一级视频免费观看 | 国产精品毛片一区二区三区 | 天堂在线观看av | 91久久精品一区二区 | 日韩av无码中文无码不卡电影 | 国产目拍亚洲精品99久久精品 | 亚洲综合伊人久久综合 | av免费一区 | 欧美精品黑人猛交高潮 | 欧美一区二区公司 | 久草一级 | 一本色道久久99精品综合 | xxxxx国产| 中文字幕在线免费观看 | 国产寡妇xxxxxxxx性开放 | 亚洲欧美在线视频观看 | 久久97久久97精品免视看秋霞 | 怡红院a∨人人爰人人爽 | 2012中文字幕在线视频 | 一本无码久本草在线中文字幕dvd | 国语精品 | 久久综合久久88 | 911亚洲精品 | 亚洲综合一区在线 | 日本精品视频免费 | 国内精品免费午夜又爽又色愉情 | 精品国产亚洲一区二区三区 | 女人扒开屁股桶爽30分钟 | 毛片最新网址 | 伊人色综合久久天天小片 | 亚洲成人在线视频播放 | 久久国产影院 | 国产精品久久久久久久久鸭 | 国产福利萌白酱在线观看视频 | 国产妇女乱一性一交 | 又湿又紧又大又爽a视频国产 | 亚洲精品乱码久久久久久按摩观 | 7777精品伊人久久久大香线蕉 | 久久精品国产欧美亚洲人人爽 | 免费又黄又爽又猛的毛片 | 久热这里只有 | 全黄色毛片 | 深夜精品| 国产超碰人人做人人爽aⅴ 国产超碰人人做人人爽av牛牛 | 精品国产乱码久久久久久移动网络 | 国产真人无遮挡作爱免费视频 | 国产美女mm131爽爽爽免费 | 尤物精品视频 | 一级久久久 | 国产一区二区三区免费看 | 夜夜嗷| 免费av观看网址 | 天天爽夜夜爽夜夜爽精品 | 久久99热狠狠色一区二区 | 国产高清区 | 四虎影视永久在线精品 | 亚洲欧美视频在线 | 能直接看的av网站 | 99国产视频 | 在线永久免费观看黄网站视频 | 免费无码又爽又刺激网站 | 136fldh导航福利微拍 | 国产女主播av在线 | 在线观看黄色国产 | 极品少妇一区二区 | 中文字幕在线有码 | 国产精品欧美一区二区三区喷水 | www久久久久久久 | 调教在线观看 | 中文字幕亚洲综合 | 亚洲猛少妇又大又xxxxx | 成人中文视频 | 日韩人妻ol丝袜av一二区 | 亚洲黄色片视频 | 国产亚洲欧美日韩亚洲中文色 | 欧美最猛性xxxⅹ丝袜 | 婷婷丁香五月天综合东京热 | 少妇高潮一区二区三区99小说 | 人妻洗澡被强公日日澡 | 国产精品三区在线观看 | 成人黄色在线视频 | 天天操天 | 婷婷久久综合九色综合88 | 天天爱天天色 | yellow免费在线观看 | 亚洲嫩 | 午夜亚洲精品 | 久久曰视频 | 九九亚洲精品 | 一区二区在线观看免费视频 | 亚洲天堂三区 | 大奶一区 | 日韩福利在线播放 | 一交一性一色一伦一区二 | 久久精品99久久久久久 | 午夜高潮视频 | 可以免费看的黄色网址 | 国产欧美日韩专区发布 | 欧美五月婷婷 | 欧洲老妇做爰xxxⅹ性视频 | 蜜臀av在线播放一区二区三区 | 一区二区三区国产 | 波多野结衣1区 | 亚洲国产精品综合久久20 | 日本xxxx自慰xxxx | 手机成人av在线 | 手机国产乱子伦精品视频 | 久久天天躁夜夜躁狠狠85麻豆 | 国产又爽又黄又无遮挡的激情视频 | 综合色站导航 | 日韩一区在线播放 | 中文字幕剧情av | 人妻互换 综合 | 黄色精品视频 | 无毒的av网站 | 丁香五月缴情综合网 | 老太婆av | 日韩人妻无码一区二区三区久久99 | 国产黄a三级三级三级老年人 | 毛片看| 老司机午夜福利av无码特黄a | 中文无码人妻有码人妻中文字幕 | 人与性动交aaaabbbb | 最新超碰 | gav成人网免费免播放器播放 | 中文字幕在线亚洲精品 | 男女黄网站 | 第四色成人网 | 国产91调教 | 国产美女精品一区二区三区 | 永久免费av | 亚洲va欧美va人人爽 | 中国色老太hd | 精品黑人一区二区三区久久 | 欧美午夜激情影院 | 午夜小网站 | 亚洲精品乱码久久久久久蜜桃91 | 久久国产精品免费视频 | 亚洲国产精品毛片av不卡在线 | 欧美日韩国产成人在线 | av日韩网址 | 亚洲欧美日韩国产成人一区 | 亚洲色图欧美日韩 | 无码人妻aⅴ一区二区三区有奶水 | 国产在线精品一区二区三区不卡 | 一个人看的www免费视频在线观看 | 草草影院在线观看视频 | 免费精品99久久国产综合精品应用 | 噜噜噜视频在线观看 | 日韩高清在线观看 | 永久av网站| 污视频网站在线 | 亚洲国产精品写真 | 人妻 色综合网站 | 欧美性猛交xxxx乱大交蜜桃 | 国产一区二区在线视频观看 | 性欧美大战久久久久久久久 | 国产日韩av免费无码一区二区三区 | 精品久久久久久无码人妻热 | 日韩精品一卡2卡3卡4卡乱码的功能 | 免费黄色网址在线观看 | 亚洲男人的天堂网站 | 日韩欧美在线综合网另类 | 中文字字幕在线 | 久久视频这里只精品 | 69精品久久 | 日本成a人片在线播放 | 久久亚洲一区二区三区明星换脸 | 图片区小说区另类春色 | a级黄色片免费看 | 成人免费无码av | 日本成人在线观看网站 | 亚洲精品久久区二区三区蜜桃臀 | 潮喷无码正在播放 |

    電子發(fā)燒友

    中國電子工程師最喜歡的網(wǎng)站

    • 2931785位工程師會員交流學(xué)習(xí)
    • 獲取您個性化的科技前沿技術(shù)信息
    • 參加活動獲取豐厚的禮品