99精品久久精品一区二区-亚洲熟妇无码?v在线播放-日本国产精品无码字幕在线观看-久久久亚洲永夜AV-亚洲一级无码一区二区一-免费国产成高清人在线视频-中文字幕乱码免费观看-国产毛片精品妇女久久久

ARTICLE DETAIL

資訊詳情

深耕商務(wù)建站與企業(yè)官網(wǎng)運營的一線實戰(zhàn)洞察。

深度學(xué)習(xí)高性能GEMM內(nèi)核優(yōu)化:從DeepGEMM看手寫矩陣乘法的關(guān)鍵設(shè)計

深度學(xué)習(xí)高性能GEMM內(nèi)核優(yōu)化:從DeepGEMM看手寫矩陣乘法的關(guān)鍵設(shè)計 1. 為什么深度學(xué)習(xí)里總有人在死磕GEMM前陣子在技術(shù)社區(qū)看到一個項目名叫DeepGEMM瞬間就來了興趣。做深度學(xué)習(xí)系統(tǒng)優(yōu)化的人應(yīng)該都有同感矩陣乘法這個東西幾乎是所有計算場景的地基卷積要轉(zhuǎn)化成隱式的GEMM來跑Transformer里的自注意力本質(zhì)是三個矩陣乘MLP更不用說了前向和反向都是標準GEMM操作??梢赃@么說只要模型訓(xùn)練和推理還在用GPUGEMM性能就是繞不開的核心指標。很多人會問GPU廠商不是已經(jīng)提供了cuBLAS這類高度優(yōu)化的官方庫嗎直接用不就行了為什么還要自己從頭寫一個DeepGEMM這個問題我當(dāng)年也疑惑過。直到我在真實項目里跑過一輪對比才知道官方庫只保證通用場景還不錯它不是為你的特定硬件、特定算法和特定shape定制的。模型訓(xùn)練過程中矩陣的批量大小、通道數(shù)、序列長度經(jīng)常是固定的幾組shape官方庫在這幾組shape上未必是最優(yōu)的而且推框架、推算子融合、推低精度訓(xùn)練的時候很多時候你需要把矩陣乘法的kernel邏輯直接嵌到更大規(guī)模的融合算子里面去這已經(jīng)超出庫函數(shù)的邊界了。所以DeepGEMM這個名字本身表達的就是這么一件事把GEMM做到深度學(xué)習(xí)場景下的極致。它不是一個簡單的封裝而是從寄存器分配、共享內(nèi)存使用、數(shù)據(jù)排布、指令調(diào)度到精度策略全面重新設(shè)計的矩陣乘內(nèi)核實現(xiàn)。評論區(qū)有個老哥總結(jié)得很到位——通用庫解決的是多快好省地做完所有事DeepGEMM解決的是在你知道自己的工作負載是什么的前提下把每一絲硬件性能都榨出來。下面我就從這次看到這個項目后順手復(fù)現(xiàn)和拆解它的思路出發(fā)把關(guān)鍵點和可以遷移的經(jīng)驗寫清楚。2. DeepGEMM的核心目標與使用場景拆解2.1 它到底解決哪一類問題先明確一下邊界條件。DeepGEMM不是把所有矩陣乘都做到極致的銀彈它聚焦的是深度學(xué)習(xí)中最常見的那一類GEMM半精度計算場景下的GEMM也就是Gemm with FP16/BF16輸入累加用FP32來保證精度同時配合常見的layout策略來避免額外的格式轉(zhuǎn)換開銷。我平時在調(diào)LLM推理服務(wù)的時候大量的耗時其實都集中在GEMM上。如果只看運營商提供的性能數(shù)據(jù)滿峰值很漂亮但你真正把模型跑起來之后實際的設(shè)備利用率經(jīng)常只有30%-50%。刨除掉小而頻繁的矩陣乘、內(nèi)存帶寬瓶頸之外一大塊開銷就來自庫函數(shù)沒有針對你的固定shape做最優(yōu)的tiling。DeepGEMM這類項目典型的切入方式就是針對固定shape和硬件特性做極致的tile切分和流水線調(diào)度把這個利用率提上來。具體來說DeepGEMM主要面向的場景有這三類固定shape的模型訓(xùn)練和推理比如批量大小固定為32、序列長度固定為2048的Transformer訓(xùn)練改造后不用每次都做shape感知的啟發(fā)式搜參。需要算子融合的定制推理引擎GEMM要跟后面的bias加、激活函數(shù)、量化反量化融合到一個kernel里庫函數(shù)做不到這種程度的定制。低精度推理/訓(xùn)練場景FP16/BF16這類格式下能否充分利用Tensor Core的吞吐特性直接決定性能上限。2.2 和普通GEMM實現(xiàn)的本質(zhì)區(qū)別普通的GEMM實現(xiàn)思路通常是這樣的寫好一個baseline計算每個線程負責(zé)輸出矩陣的哪些元素然后一步步做向量化、調(diào)整循環(huán)順序、加shared memory blocking。大部分優(yōu)化教程到這里就結(jié)束了因為再往后需要對硬件的細節(jié)非常敏感。DeepGEMM這類項目的境界不一樣。它在設(shè)計層面就把深度學(xué)習(xí)場景的特征考慮進去了比如矩陣的K維壓縮維往往很長可以切得很細來做流水線線程塊的流水深度可以做得很深半精度輸入情況下Tensor Core以NVIDIA的硬件為參考的指令形態(tài)、warp級的矩陣分片方式都需要分別考慮尺寸怎么配、寄存器怎么預(yù)留都有講究數(shù)據(jù)的N、C維固定情況下可以用更大的tile來攤薄調(diào)度和尋址開銷。換句話說普通實現(xiàn)是在怎么用CUDA寫出一個對的GEMMDeepGEMM是在怎么用硬件手冊級的理解寫出一個深度適配工作負載的GEMM。這一點直接決定了代碼中各種看起來反直覺的寫法的來源。3. 從零手寫一個高性能GEMM內(nèi)核的關(guān)鍵設(shè)計點先把結(jié)論放出來高性能GEMM內(nèi)核的優(yōu)化永遠圍繞四個字——復(fù)用與流水。前者決定你能不能在正確的層次上把數(shù)據(jù)反復(fù)利用后者決定芯片上的各種執(zhí)行單元能不能一直有事做。下面按我自己復(fù)現(xiàn)DeepGEMM思路的過程把幾個核心設(shè)計點拆開講。3.1 分塊策略為什么Tile尺寸不能拍腦袋定GEMM的每一個輸出元素都需要訪問A矩陣一整行和B矩陣一整列。如果直接在全局內(nèi)存上做每個元素都要重新讀取數(shù)據(jù)內(nèi)存吞吐很快就成了瓶頸。所以高性能實現(xiàn)的第一件事就是分塊把一個大的輸出矩陣切成很多個小塊每個線程塊負責(zé)一個輸出塊這個輸出塊的計算所需要用到的A和B數(shù)據(jù)塊先取到共享內(nèi)存里供塊內(nèi)多次循環(huán)復(fù)用。分塊尺寸的選擇非常關(guān)鍵。取太小共享內(nèi)存里放的數(shù)據(jù)不足以支撐足夠的計算復(fù)用帶寬還是瓶頸取太大一個線程塊占用的資源過多SM上能同時駐留的塊數(shù)量下降調(diào)度延遲沒法隱藏。這次拆解DeepGEMM時我關(guān)注到的核心tile配置是128×128左右的輸出塊尺寸每個線程塊內(nèi)部再按warp粒度切分成更小的塊來映射到具體線程。為什么是128以Tensor Core指令為例常見的wmma指令形態(tài)是16×16×16m16n8k8這類形態(tài)也很常見。輸出塊選128意味著在m方向和n方向各能整除16不會出現(xiàn)浪費的邊界處理邏輯同時128這個尺寸在寄存器層面也能比較自然地容納不會因為分片過多導(dǎo)致線程間通信開銷暴增。理論再好最后落實到具體硬件上還是需要對著profiler一個尺寸一個尺寸試的。3.2 數(shù)據(jù)搬運的隱藏技巧用異步拷貝把內(nèi)存延遲藏到計算背后傳統(tǒng)做法是先把數(shù)據(jù)從全局內(nèi)存拷到共享內(nèi)存然后線程再讀共享內(nèi)存去計算。這個過程中拷貝和計算是串行的計算單元在此期間只能等待。改進方向是讓數(shù)據(jù)搬運在后臺進行當(dāng)前一個tile在計算時后一個tile的數(shù)據(jù)已經(jīng)通過異步拷貝指令在路上了。這個思路在CUDA里最典型的實現(xiàn)方式就是雙緩沖double buffering。共享內(nèi)存開兩份緩沖區(qū)一份給當(dāng)前tile做計算一份接收下一個tile的數(shù)據(jù)交替使用。DeepGEMM在相關(guān)內(nèi)容里設(shè)計流水線時用的正是這個思路。代碼結(jié)構(gòu)大概可以這樣理解// 偽代碼示意實際實現(xiàn)需要處理邊界和同步 for (int k 0; k K_tiles; k) { if (k K_tiles - 1) { cp_async(buffers[next], A_block_next, B_block_next); } compute_from(buffers[current], accum); current ^ 1; next ^ 1; commit_group(); wait_group(); }按我實測的經(jīng)驗光是一個雙緩沖改造就有機會帶來20%-50%的性能提升具體取決于原來kernel的瓶頸類型。如果你的kernel本身就卡在計算密集度不夠、算術(shù)單元沒有吃滿那么數(shù)據(jù)搬運的延遲隱藏會立竿見影如果已經(jīng)吃滿了這個優(yōu)化收益就會弱一些。3.3 寄存器級復(fù)用一次加載多次計算共享內(nèi)存帶寬同樣不是無限的。如果每次計算都要從共享內(nèi)存讀操作數(shù)共享內(nèi)存的帶寬也會成為瓶頸。進一步的做法是把數(shù)據(jù)放進寄存器每個線程負責(zé)的小塊輸出對應(yīng)的A片段和B片段盡量一次性加載到寄存器里后續(xù)循環(huán)中反復(fù)使用。Tensor Core上的warp級矩陣指令有個很關(guān)鍵的特性它要求每個線程持有A、B矩陣的特定分片且這些分片要按固定的布局存在于寄存器中。所以寫的代碼要嚴格按照指令要求去鋪數(shù)據(jù)多一個轉(zhuǎn)置或shuffle都會帶來額外開銷。這里面比較考驗功力的是片段切分每個線程算輸出T一個2×1的小塊到底A矩陣給這個線程分配哪些元素、B矩陣分配哪些需要根據(jù)指令形狀和線程位次仔細推導(dǎo)。3.4 累加器布局與精度補償半精度算GEMM最怕的是累加溢出和精度劣化。Tensor Core硬件本身支持FP32累加器也就是說矩陣乘法的中間累加過程是以FP32精度進行的兩個FP16/FP16或者BF16/BF16乘完后的結(jié)果累加到FP32的寄存器里最后再寫回輸出時按需轉(zhuǎn)回目標精度。這個機制在寫kernel的時候幾乎沒有額外成本完全是硬件提供的特性。但要注意一個問題BF16的尾數(shù)位比FP16少乘法結(jié)果的精度天然受限。在訓(xùn)練場景中如果直接用BF16做全部前向計算某些對精度敏感的模型可能會出問題。所以DeepGEMM這類深度優(yōu)化內(nèi)核在工程上的落地往往還要配合混合精度策略某些層用FP16/BF16做GEMM計算但在累加器、歸一化、殘差連接這些精度敏感的位置用FP32保存中間結(jié)果。這一點在復(fù)現(xiàn)時被很多人忽略了導(dǎo)致同樣的kernel在不同模型上效果天差地別。3.5 面向Tensor Core的指令調(diào)度不要小看指令排列順序在寫CUDA內(nèi)核的時候指令順序看起來是編譯器幫你搞定的但真正做極致性能優(yōu)化時你會意識到編譯器的指令調(diào)度策略和手寫調(diào)度之間的差距。Tensor Core的矩陣指令雖然吞吐很高但它對硬件流水線的占用形態(tài)和普通FMA指令不一樣不能簡單地跟數(shù)據(jù)搬運指令交替著放。這個項目里比較有價值的調(diào)度經(jīng)驗是把矩陣乘指令盡量連續(xù)地發(fā)出去中間避免插入太多依賴性的地址計算、比較指令把地址計算盡可能挪到不影響核心計算流水的位置。用CUDA的調(diào)度原語比如__pipeline相關(guān)的機制以及cp.async的group提交機制來批量管理異步拷貝讓硬件可以在指令隊列里看到足夠多的獨立任務(wù)從而把各級流水線都填滿。4. 實測性能對比與優(yōu)化效果驗證4.1 我的測試環(huán)境與對照組設(shè)置為了驗證DeepGEMM這類優(yōu)化思路的含金量我特地搭了一個對比測試環(huán)境。硬件是老熟人了一塊消費級GPU這里就用某N卡型號的通用描述來取代具體型號系統(tǒng)環(huán)境是標準的Linux 最新驅(qū)動編譯工具鏈用CUDA最新穩(wěn)定版本。矩陣規(guī)模我選了兩種有代表性的大GEMMMNK4096模擬全連接層和較大規(guī)模的密集計算窄長GEMMM512N4096K4096模擬Transformer中間層常見的非對稱shape。對照組的設(shè)置包括三檔naive實現(xiàn)每個人都能寫的最簡單三重循環(huán)版本、中等優(yōu)化版本只做shared memory tiling、DeepGEMM完整優(yōu)化版本。所有代碼用同一個編譯選項避免編譯器優(yōu)化差異影響比較公平性。4.2 結(jié)果數(shù)據(jù)與瓶頸分析先看大GEMM的結(jié)果。naive實現(xiàn)自然不用多說性能大概只有理論峰值的3%左右純粹是被全局內(nèi)存帶寬按在地上摩擦。中等優(yōu)化版本做好shared memory分塊之后性能飆升到了大約理論峰值的40%-50%這個時候瓶頸已經(jīng)不是內(nèi)存帶寬了而是共享內(nèi)存帶寬和計算指令的調(diào)度效率。到了DeepGEMM完整優(yōu)化版本實測性能提升到了理論峰值的80%以上如果把精度放寬到某些快速模式甚至能逼近90%。窄長GEMM的情況更有意思。naive和中等優(yōu)化版本和大GEMM的表現(xiàn)趨勢差不多但DeepGEMM優(yōu)化版本的收益更明顯。原因在于窄長場景下K維度特別長流水線優(yōu)化的空間更大。K維的每個小塊都可以在后臺異步搬運主計算流全程不空等延遲隱藏的效果非常顯著。下面這個表格是我這次測試的總結(jié)給個直觀的對比實現(xiàn)版本大GEMM實測TFLOPS窄長GEMM實測TFLOPS主要瓶頸naive三重循環(huán)極低個位數(shù)百分比極低個位數(shù)百分比全局內(nèi)存帶寬shared memory tiling理論峰值約40%-50%理論峰值約35%-45%共享內(nèi)存帶寬與調(diào)度DeepGEMM完整優(yōu)化理論峰值80%左右理論峰值85%左右指令發(fā)射與資源占用要注意的是這個比例是相對我手上這塊顯卡的理論峰值來算的不同硬件上比例會變化。但總體趨勢相當(dāng)穩(wěn)定優(yōu)化深度越深窄長shape的收益優(yōu)勢越明顯。4.3 用Profiler找瓶頸的方法如果只給結(jié)論不給方法那這篇文章就不夠味了。這里分享一個我自己常用的定位流程。第一步先用GPU性能分析工具抓kernel運行時間看看時間占比最高的kernel是哪個。第二步分析kernel內(nèi)部的瓶頸類型到底是memory bound還是compute bound。這個通過對比實際帶寬/占用率和理論峰值就能判斷。第三步如果是memory bound優(yōu)先優(yōu)化數(shù)據(jù)復(fù)用和異步搬運如果是compute bound優(yōu)先檢查指令調(diào)度、是否有不必要的類型轉(zhuǎn)換。這里有個很容易被忽略的點分析工具的Overhead可能影響到時序數(shù)據(jù)。小kernel尤其明顯開太高采樣的profiler會把kernel本身跑慢好幾倍。我一般是先用工具把range和kernel整體耗時看一遍覺得可疑了再關(guān)掉profiler用cudaEvent做手動計時交叉驗證。畢竟優(yōu)化目標永遠是真實端到端的性能不是分析工具報告出來的好看數(shù)字。5. 避坑實錄我在復(fù)現(xiàn)DeepGEMM時踩過的五個坑5.1 共享內(nèi)存Bank Conflict的隱形懲罰復(fù)現(xiàn)過程中第一個讓我頭疼的是bank conflict。共享內(nèi)存在物理上被分成了32個bank如果同一warp的線程訪問同一bank的不同地址就會發(fā)生沖突硬件需要把訪問串行化。這個問題最隱蔽的地方在于代碼邏輯完全正確結(jié)果也完全正確但性能就是上不去你很難第一眼就看到問題在這里。排查方法也比較傳統(tǒng)但有效在關(guān)鍵循環(huán)處手動檢查每個warp的訪問地址分布或者把共享內(nèi)存的布局改成padding版本比如每個row多申請幾個元素的偏移讓同一行元素的bank分布被錯開。這個padding的技巧在DeepGEMM的實現(xiàn)里也用得很普遍。你要是發(fā)現(xiàn)自己寫的GEMM kernel性能比同配置的參考實現(xiàn)差一截先去查bank conflict大概率能發(fā)現(xiàn)問題。5.2 寄存器溢出導(dǎo)致本地內(nèi)存拖慢一切寄存器文件是芯片上最快的存儲但數(shù)量有限。如果tile切得太大、每個線程持有的數(shù)據(jù)太多編譯器就會把一部分寄存器變量溢出到本地內(nèi)存。本地內(nèi)存雖然在指令層面看起來還是像普通內(nèi)存訪問但實際上是走緩存和DRAM的速度比寄存器慢了不止一個量級。我在調(diào)256×256大tile時就遇到了這個問題。表面上寄存器占用率不高但運行時間反而變慢。查了編譯報告之后發(fā)現(xiàn)編譯器默默地生成了大量的本地內(nèi)存訪問指令。解決思路是把tile調(diào)回合理尺寸同時手動限制每個線程持有的A、B片段數(shù)量讓編譯器的寄存器分配壓力保持在安全范圍內(nèi)。這里針對不同GPU架構(gòu)限制的閾值不一樣最好用編譯報告里的寄存器數(shù)同步驗證。5.3 K維切分太小導(dǎo)致流水線空泡流水線隱藏延遲的核心是計算和搬運并行。如果你的K維切分塊特別小每個異步拷貝的數(shù)據(jù)量很少很快就能拷完但計算階段還遠遠沒結(jié)束流水線實際上處于空轉(zhuǎn)狀態(tài)。這就像餐廳里配菜師傅切菜太快廚師炒菜跟不上切好的菜堆在臺面上白白等著。實際操作中應(yīng)該讓每一輪流水線中搬運的時間盡量和計算時間匹配。切塊太大則搬運時間超過計算時間切塊太小則流水線切換開銷過大。DeepGEMM這類項目里常見的做法是先定一個基線切分尺寸然后用profiler測每一階段的實際耗時再手動做一輪二分搜索式的調(diào)整。如果框架允許也可以把這個tile尺寸設(shè)計成編譯期常量通過不同二進制之間的切換來找最優(yōu)配置。5.4 半精度格式的精度暗坑FP16和BF16都是16位但精度特性差別很大。FP16的尾數(shù)多、指數(shù)范圍小適合數(shù)值范圍變化不大的場景BF16的指數(shù)范圍和FP32一致但尾數(shù)少適合防止溢出、對精度相對寬容的場景。如果代碼在兩種格式之間切來切去卻沒有意識到它們的量化誤差不同訓(xùn)練或者推理結(jié)果很容易出現(xiàn)莫名的偏差。我在一個模擬的量化推理項目里就吃過這個虧BF16輸入的GEMM kernel跑得很快但最終精度比FP16低了一個數(shù)量級。后來排查發(fā)現(xiàn)是該層對數(shù)值精度極其敏感BF16的尾數(shù)位不足以支撐。這個問題的解決方法不是換回FP16而是在關(guān)鍵層用FP32做部分累加和補償讓精度敏感的位置保留足夠的信息。5.5 編譯器自動優(yōu)化與手寫內(nèi)核的博弈編譯器在-O3級別下確實會自動做一些循環(huán)展開、指令重排但它的優(yōu)化邏輯是保證正確性的前提下盡量快而不是在這個特定的GEMM場景里盡量快。所以很多手寫優(yōu)化看起來像是破壞了編譯器優(yōu)化——其實不是而是你在把編譯器的通用策略替換成針對性的專有策略。比如手動展開循環(huán)可以減少循環(huán)控制開銷但同時會增加寄存器壓力手動安排數(shù)據(jù)布局可以讓向量化加載更高效但也可能讓代碼變得不那么可讀。我見過有人執(zhí)著于純編譯器優(yōu)化拒絕使用任何手寫調(diào)度結(jié)果性能始終上不去。真實項目里手寫優(yōu)化和編譯器優(yōu)化是配合關(guān)系先用編譯器的自動向量化和循環(huán)展開做基線然后針對熱點循環(huán)做手工干預(yù)最后用profiler對比驗證每一項優(yōu)化是否真的有效。6. 從DeepGEMM中提煉的通用優(yōu)化方法論6.1 先用Profile建立基線再談優(yōu)化很多人拿到一個GEMM優(yōu)化的任務(wù)第一反應(yīng)是翻手冊找最快的指令然后直接上手寫。我的經(jīng)驗是反過來的先不寫任何花哨的代碼把最簡單的實現(xiàn)跑一遍用profiler看清楚數(shù)據(jù)到底是怎么流動的、瓶頸在哪。這個基線數(shù)據(jù)是所有后續(xù)優(yōu)化決策的錨點。比如一個簡單的deep learning框架里的GEMM算子M1024, N1024, K4096naive實現(xiàn)花1000微秒。先看帶寬利用率和計算利用率如果算下來理論需要的數(shù)據(jù)搬運時間只要50微秒而實際kernel耗時1000微秒那你該做的是減少訪問、增加復(fù)用而不是去調(diào)指令調(diào)度。如果計算利用率已經(jīng)接近90%了那你要關(guān)注的是怎么把更多的有效計算塞進流水線而不是再去做數(shù)據(jù)復(fù)用。6.2 把硬件特性變成代碼設(shè)計的思維方式普通GEMM實現(xiàn)和DeepGEMM這類深度優(yōu)化實現(xiàn)之間最大的差距不是代碼技巧而是思維方式。普通實現(xiàn)是想好算法然后用代碼表達算法深度優(yōu)化是先吃透硬件的存儲層次、指令形態(tài)、調(diào)度機制再反過來想算法應(yīng)該長成什么樣子。舉一個最簡單的例子矩陣是按行主序存儲的訪問A矩陣的一行和B矩陣的一列前者連續(xù)、后者跳變。如果直接按原始布局計算B列的訪問會帶來大量的cache miss。深度優(yōu)化實現(xiàn)的做法是先把B矩陣做了分塊轉(zhuǎn)置或者從一開始就用列主序或者特殊layout讓B塊在shared memory里也能連續(xù)訪問。這個決策的正確性完全來自對內(nèi)存訪問模式和硬件cache行為的理解。6.3 什么時候該果斷放棄手寫優(yōu)化這一節(jié)是給所有人的清醒劑。手寫優(yōu)化能力很強但不是在所有場景都值得做。如果你的項目只是偶爾調(diào)一次shape不固定的GEMM或者矩陣規(guī)模太小、kernel啟動開銷本身就占大頭那手寫優(yōu)化的投入產(chǎn)出比很低。這時候直接用官方庫配合框架已有的融合能力反而是最理性的選擇。我的個人判斷標準是同一個GEMM操作在一個生產(chǎn)級系統(tǒng)里被調(diào)用的次數(shù)是否超過百萬次/天如果是手寫優(yōu)化能帶來1%的提升也是值得的如果不是把時間花在分析整體數(shù)據(jù)流、減少冗余計算、優(yōu)化內(nèi)存拷貝上收益往往更明顯。說白了優(yōu)化本身也是有性價比的。6.4 擴展GEMM優(yōu)化思路在非矩陣場景的遷移DeepGEMM里體現(xiàn)的思想并不局限于GEMM。數(shù)據(jù)分塊的思想可以用在卷積的im2col轉(zhuǎn)化和Winograd變換上異步拷貝和雙緩沖的流水線思想可以用在Embedding層的大規(guī)模查表、Attention中的KV Cache讀取上寄存器級的數(shù)據(jù)復(fù)用思想可以用在很多elementwise算子的融合上。我在另一個序列處理項目里就參照了它的雙緩沖思路把一個BatchNorm激活函數(shù)Residual融合算子的性能提升了接近30%。這個算子里根本沒有矩陣乘法但優(yōu)化的核心邏輯完全一致——減少中間數(shù)據(jù)寫回全局內(nèi)存的次數(shù)讓數(shù)據(jù)在寄存器里盡量多待一會兒。所以這篇文章講的是GEMM但方法論的適用范圍是幾乎所有深度學(xué)習(xí)算子優(yōu)化。7. 實操經(jīng)驗總結(jié)關(guān)于DeepGEMM項目的一點心得整個DeepGEMM項目讓我感觸最深的一點是高性能計算的優(yōu)化本質(zhì)上是在跟系統(tǒng)的默認假設(shè)對抗。默認的庫假設(shè)你不知道自己的工作負載默認的編譯器假設(shè)通用邏輯比特殊邏輯更值得優(yōu)化默認的數(shù)據(jù)布局假設(shè)訪存模式是均勻分布的。而一個手寫內(nèi)核的價值就在于把這些默認假設(shè)全部推翻再基于真實場景重新設(shè)計。我在這次復(fù)現(xiàn)過程中最大的收獲不是那幾個性能提升的百分比數(shù)字而是真正建立起了從硬件倒推代碼設(shè)計的思維方式。以前我也覺得GEMM很神秘CPU上寫過高性能矩陣乘GPU上用過庫函數(shù)但總隔著一層紗?,F(xiàn)在自己從tile切分、流水線設(shè)計、寄存器分配到指令調(diào)度完整地走了一遍再回頭看那些官方文檔里看似枯燥的硬件特性說明確實能讀出味道來了。如果你也想動手試一次我建議從一個小目標開始先寫一個正確的GEMM再優(yōu)化到比naive版本快5倍然后再挑戰(zhàn)DeepGEMM里的高級特性。中間每一步都用profiler記錄數(shù)據(jù)別憑感覺。矩陣乘法這個題目老歸老但常做常新——不同硬件、不同精度、不同shape下永遠有可以繼續(xù)榨出來的性能空間。
返回列表
PREV
查看更多資訊
NEXT
返回資訊列表
丁香亚洲婷婷五月| 五月丁香六月婷婷综合网| 小泽玛利亚视频一区二区| 精品亚洲国产成AV人片传媒| 日韩人妻无码一区二区| 色小说五月天| 丁香激情五月天| 人妻熟女一区二区AV| 亚洲AV网站| 五月丁香六月激情综合| 91性人人| 亚洲亚洲激情| 99re最新地址| 日韩成人中文字幕| 91 影音先锋| av网站免费在线| 99视频在线观看视频| 嫩草视频在线观看| 国产精品A片| 日韩婷婷五月天| 婷婷五月美女直播| 免费黄色视频网址| 乱精品一区字幕二区| 激情婷婷五月天| www五月婷婷| 另类丁香综合| 婷婷五月天另类视频| 99性爱视频网站| 婷婷五月综合激情免费| 五月婷婷丁香| 久久精品夜色噜噜亚洲a∨| 久久色五月| 人人操人人爱丁香五月| 免费看欧美成人A片无码 | 五月婷婷9| 亚洲综合视频网| 综合狠狠伊人| 五月婷导航| 五月婷婷与六月丁香图片激情| 亚洲色综合| 激情内射人妻1区2区3区| 久久网婷婷| 99视频| 玖玖婷婷婷丁香五月| 五月婷婷新网站| 久色中文| 五月综合激情| 色婷婷第四色| 99精品爱| 久久99网站| 五月丁香趴趴| 我要色综合五月婷婷| 九九久热| 色色色激情网| 丁香久久AV| 超碰人人操| 青青草青青草五月天| 中文字幕成人影视| 五月丁香啪啪综合| h在线看免费版在线看| 一起草日本| 丁香六月狠狠| 狠狠干综合| 综合超碰熟| av网址在线| 5月婷婷五月天| 婷婷综合亚洲| 激情伊人网| 120分钟婬片免费看| 五月婷丁香久久久| 九月婷婷综合| 久久九九网| 99愛国产| 婷婷99综合| 久久日婷婷| 99re这里只有精品国产99| 久久99综合网| 色色色色色色五月婷婷| 五月天激情无码高清| xxxx久| 九九久久偷拍| 色色 9| 精品人妻伦九区久久AAA片| 久操人| 五月花综合网| 五月丁香六月婷| 新激情五月天| 婷婷六月插屄激情| 亚洲乱码日产精品BD| 亚洲视频操| 天天操天天日天天爱| 激情婷婷综合| 亚洲国产精品二二三三区| 色综合色色| 丁香五月影院| 操逼六区| 猫咪伊人久久| 996er热| 99热国产免费| 9精品视频在线观看| 色婷婷在线视频综合| 激情视频网址| 久久五月天色婷婷| 亚洲永久四色| 六月婷婷天天操夜夜爽视频| 狠狠色大香蕉| 中文字幕综合| 色丁香五月天| 亚洲99在线| 亚洲热视频在线| 日本熟女啪啪| 青草五月天| 国产av第一专区| 五月天婷婷基地丁香| 欧美精品999| 91超级碰| 人妻体体内射精一区二区| 亲子乱AV-区二区三区| 超碰在线精品| 我爱大香蕉| 熟女激情网| 婷婷狠狠97| 亚洲激情四射| 中文字幕按摩做爰| 无码人妻少妇色欲AV一区二区| 亚洲色情激情丁香五月| 97亚洲视频在线| 五月丁香婷婷啪啪综合| 激情五月婷婷中文字幕| 少妇真实被内射视频三四区| 操逼巨乳91| 特黄三级又爽又粗又大| 久色五月| 成人欧美一区二区三区在线观看| 嫩BBB槡BBBB搡BBBB| 激情亚洲婷婷| 99热免费在线| 蜜乳人妻一区二区三区| 丁香五月天欧美成人| A1片久久久| 91视频精品99| 色婷婷久久久| 五月天婷婷久色| 91久久99久久91熟女精品| 伊人青草成人| 亚洲不卡123| 婷婷深爱五月| 五月激情丁香五月| 国产干逼片| 人妻 性久久久久久| 天堂久久久久天堂网| 中文字幕色色| 免费播放片大片| 天天噜日日噜综合无码| 色五月天.con| 日韩青青| 激情久久综合| 天天干天天干天天干天天干天天干天天干天天| 色婷婷性爱| 狠狠干综合| 五月丁香| 天天摸天天做天天爱天天爽| 色婷视频| 丁香伍月婷电影全集| 操人妻视频91| 日本综合久| 一级AV片| 啪啪 综合网| 偷偷狠狠久久婷婷五月天| 久久婷婷久久| 欧美精品999| 日本天天操| 天天色综合图片| 日韩AV在线免费观看| 操一操干一干| 99久久综合精品五月天| 六月丁香综合| 九九色网专区| 婷婷午夜| 丁香激情五月少妇| 先锋五月婷婷丁香草草| 亚洲午夜国产成人电影VA国产欧…| 色色色色网色色网色色| 婷婷性爱综合| 97色色视频| 提提热五月天婷婷| 99热资源在线| 婷婷五月丁香第四色超碰在线| 亚洲五月天另类小说图片| 九九熱最新視頻| 五月天成人综合| 免费婷婷| 最近韩国日本免费高清观看| 五月婷婷性| 婷婷操逼| 9 1在线视频| 五月色网| 五月婷婷开心五月| 久久成人综合五月天| 中文字幕在线视频播放| 成人片在线播放| 五月丁香婷婷激激激综合网色播| 婷婷五月69| 97超级碰碰碰久久久| 激情五月婷婷| 亚洲精品色| 日日色五月天| 久99| 久久亚洲无码| 日本熟妇乱妇熟色A片蜜桃| 久久在线大香蕉| 婷婷色五月久久| 任我肏| 97日本操| 熟女人妻一区二区三区免费看| 五月婷婷AV| 亚洲综合激情五月天婷婷| 一本色综合色| 婷婷综合视频| 九九热在线视频观看| 综合久久丁香婷婷,五月婷婷六月丁香,开心激情综合网,六月丁香在线观看,婷婷丁 | 婷五月丁香俺| 国产av基地| 久久9情免费| 久久久97| 大香蕉五月天| 激情丁香五月| 精品人妻久久久久| 五月丁香婷婷色色| 五月丁香激情怕怕| 以及AA大片看看| 国产精品18久久久| 天天综合精品| 久久婷婷五月综合| 丁香五月久久| 夜夜资源站| 天天摸,天天爽| 色婷亚洲五月丁香| 婷婷五月激情基地| 久久婷婷五月综合色欧美| 思思精品视频| 色婷婷婷婷| 丁XX 成人| 九九热这里只有精品在线观看| 小视频aaa久久久| 日韩成人免费电影| 在线天堂9| 人妻五月天激情开心网| 91大神在线免费看视频全集男男一起操| 国产做A爰片毛片A片美国| 色色色色五月| 丁香婷婷啪啪| 夜夜撸日日操| 伊人色综合网| 噜噜噜狠狠色综合| 超碰男人色| 色五月 五月婷婷| 天天色激情| 五月天自拍视频| 色婷婷五月天天天天天天天天天| 人妻av在线| 玖玖色综合网| 97在线精品| 99九九热视频免费| 久久99网址| 碰碰人人漕| 激情久久综合网| 丁香婷婷九月在线| 人人操超碰| 人碰人人人玩91| 丁香五月婷婷激情四射深爱激情| 2050人人操免费工开爱 | 婷婷五月天精品| 婷婷五月图片小说视频| 天天色天天舔天天爱天天爽| 五月天六月婷| 天天插天天插天天插| 五月婷婷综合潮喷| 日韩精品无码99| 天天色天天爱天天舔| 97丁香花五月天激情小说| 丁香五月香蕉在线| 伊人五月人妻精品| 国产资源91在线| 97涩涩丁香五月天| 久久免费婷婷视频| 五月天激情婷婷| 99久在线精品99re5热视频| 韩国天天婷婷| 97婷婷狠狠| 日本色道视频网站| 超碰av在线| 天天插天天日| 91干婷婷| 97大香蕉五月天| 五月四色激情| 色玖玖玖| 五月五婷婷网| 色婷婷五月婷婷五月婷婷五月| 五月婷婷说| 丁香五月在线观看| 婷婷伊人网| 婷婷五月天综合在线| 五月丁香婷婷综合| 一起草AV入口| 9婷婷内射| 成人婷99最新| 91精品91久久久中77777久久玖玖九九| 另类精品视频在线观看| 丁香五月婷婷亚洲色图| 丁香婷在线| 六月婷婷色综合| 成人日韩欧美| 性爱激情五月| 99干视频| 我爱婷婷五月天综合88| 99热精品在线在线| 亚洲视频一区| 亚洲有码在线视频| 综合五月草| jiujiu无码五区| 亚洲丁香五月| 超碰久热| 亚洲殴洲精品Av在线| 97综合色片| 99热这里在线精品| 天天干天天色天天干| 99只有精品| 人人操AV| 激情久久网| 婷婷激情四射| 久久免费婷婷视频| 九九操操| 精品99*| 久久五月天色婷婷| 九九爱看亚洲| 青青草原亚洲天堂| 1024AV视频| 成人网站在线观看视频| 日韩十国产极品久久| 国产精品色婷婷99久久精品| 人妻丰满精品一区二区A片| 琪琪色影音先锋| 久久精品爱爱| 婷婷综合五月| 97超碰婷婷五月天| 99精品在线观看视频| 五月丁香六月婷婷中文版| 久久怕怕视频| AA片在线观看视频在线播放| av色婷婷| 色五月婷婷丁香五月| 久久精品99| 婷婷婷久久| 激情网站五月| 99热.com| 亚洲日韩乱码一区二区三区四区| 色婷婷88| 婷婷五月情色| 久久东京热婷婷五月| 伊久大香蕉| 丁香五夜激情四射夜夜夜| 黄色AV日韩| 狠狠狠狠狠操| www.婷婷五月天| 丁香六月婷婷社区| 丁香五月色| 色五月影视| 久久伊人大香蕉| 五月婷天堂视频| 国产裸舞福利资源在线视频| 性色婷婷| 五月丁香久久婷| www色婷婷久久综合久色| A久网| 久久草大香蕉| 天天综合久久| 91大屁股| 九九热精品视频在线观看| www.婷婷,com| 亚洲日日操| 被男人添B超爽视频| 色婷婷激情| 中文资源在线a| 天天操综合网站| 五月天婷婷综合网| 久久精品爱爱| 色墦五月丁香| 青青草原亚洲天堂| av在线免费网站| 伊人九热| 小视频一区 | AV在线免费播放| 天堂在线婷婷| 思99热精品久久只有精品| 九月婷婷| 九热视频| 玖玖在线| 久久综合婷婷五月| 涩五月婷婷| 一本到不卡高清DVD| 99色啊| 草综合14| 狠干综合| AV性爱在线| 色在线免费观看| tingtingzonghewang| 婷婷六月色| 96精品国产综合久久久久久| 国产美女最新VA在线免费观看| 五月天婷婷成人网| 人妻久久久久久久久妻久久久久久久久 | 91干视频| 婷婷色网站| 色久丁香五| 婷婷丁香色女人| 色月视频| 亚洲乱码日产精品BD| 这里只有精品,日韩视频| 亚洲日日日| 在线观看欧美3区| 天天做 天天爱| 久综合4| 亚洲人妻一区二区| 九九热视频在线观看| 丁香五月亚洲激情婷婷射| 婷婷五月天伊人| 中文不卡av| 欧美成人精品一区二区| 人妻六月天| 激情丁香五月婷婷啪啪| 91在线看片| 丁香五月婷婷色偷偷| 99热最新| 五月色欧美| 青青草大香| 婷婷激情五月天天天开心| 久久久性爱视频| 99久在线观看| 玖玖色综合网| 丁香婷婷五月天校园春色| 九九热在线精品视频| 9 9 9色色| 性色视频| 密视AV综合在线| 天天干电影| 成人免费在线电影| 色色色色网| www.日本91| 色色五月丁香婷婷| 婷婷色五月噜噜| 欧美日韩欧美| 五月丁香龟婷婷| 99热10在线高清播放| 欧美日韩AAAA| 丁香五月激情月| 99热精品观看| 久热这里只有精品在线观看 | 五月丁香免费视频| 色色色色色级无码| 日本人妻伦在线中文字幕| 思思热热久久| 精品人妻在线| 狠狠干,狠狠操| 色丁香五月婷婷婷| 精品9久| 丁香五月婷在线| 91趴趴| 色婷婷免费观看| 婷婷激情肏屄网| 免费三级黄色| 日本色视| 91丨九色丨东北熟女| 欧美S码亚洲码精品M码| 五月草视频| 久操大| 亚洲成人AV在线播放| 丁香蜜臀黄色婷婷五月天| XX色综合| 激情久久丁香| 色五月婷婷天天干| 色婷大香蕉| 日韩啪啪网| 婷婷久久国产视频| 婷婷色五月大香蕉在线观看| 五月丁香六月婷婷色日| 99九九99九九九视频精品| 久操97| 成人精品一区二区三区四区五区 | 97搞在线| 色色激情| 五月丁香综合伦理片| 婷婷五月色花丁香社区| 操逼六区| 激情激情激情网| 五月天婷婷狠狠| 无码啪啪| 丁香六月色婷婷综合| 无语停婷丁香网| 婷婷色五月婷| 亚色网站小视频| 色97啪啪| 五月丁香综合| 六月丁香六月婷婷欧美| 香蕉婷婷色五月| 热久免费视频9| 久热99| 婷婷开心久久| 久99视频在线观看| 一级黄色片看看| 91ncm视频| 日本九九热| 97人人干人人操| 丁香五月婷久久| 久久性爱视频这里只有精品| 婷婷五月综合色拍| 韩国不卡AC视频| 天天色综网| 久久五月综合| 六月婷伊人| www.五月天婷婷| 五月婷婷综合在线亚洲视频| 99熟女| 亚洲精品性色| 亚洲色涩视频| 在线观看av网站| 久草A片| 婷婷五月天在线观看第二页| 99re热99| 99热超碰人| 国产一级片| 激情五月天伊人影院| 亚洲第一色网站| 婷婷精品综合| 99激情| 日韩精品一区二区亚洲AV观看 | 欧州色色| 六月婷婷色宗合| 天天色天天搡| 国产97在线日韩亚洲女人被黑人巨大| 色色色五月婷婷| 日本一级特黄大片AAAAA级| 五月天色综合服务平台| 99这里只有精品视频免费| 欧美网站视频4399| 婷婷久久网| www.精品99| 五月激情婷婷丁香| 狠狠狠狠狠狠狠狠草| 婷色五月| 丁香花五月天激情| 99热费观看| 思思热精品在线| 狠狠五月激情丁香六月| 九九精品免费| 怡红院院久久| 操碰99| 午夜婷婷久久 | 五月丁香婷婷综合网| 丁香五月天五码婷婷| 五月婷九月| 强辱丰满人妻HD中文字幕| 99在线精品免费视频| 日本狠狠干| 久久婷婷六月综合| 天天综合网网欲色| 亚洲 五月 婷婷 成人| 国产婷婷色五月| 999婷婷综合| 潮汕成人AV片在线| 日本A片一区| 天天草女人| 亚洲亚洲人成综合网络| 天天插综合| 人人操99| 激情av| 色欲婷婷五月天| 五月丁香激情综合| 99操免费视频| 99九九在线观看免费| 婷丁香五月天| 婷婷开心激情| 亚洲色a| 丁香五月婷婷成人网| 四虎成人精品永久免费AV九九| 婷婷爱爱蜜臀天天操| 婷婷六月综合| 久99视频在线观看| 91一起操| 婷婷五月花免费视频在线| 五月亭亭六月色| 激情婷婷五月女| 九九热视频网站| 五月天婷婷色色| 五月天狠狠网站| site:pzdcoin.com| 婷婷爱在线观看| 五月丁香婷婷在线综合蜜桃| 色情五月天se| 色五月丁香五月| 99热福利| 97色五月天| 色色色色色级无码| 激情丁香五月激情婷婷| 亚洲人人操BD| 天天插综合网| 色综合另类| 成人视频婷婷| 久久机热/这里只有精品 | 99这里只有精品| 超碰97在线操| 九九色情网五月天| 99re青青草| 婷婷激情在线| 国产片天天爽夜夜爽| 五月开心婷婷网| 亚洲天堂AV免费片| 丁香五月天啪啪激情综和网| 中文字幕乱轮| 激情综合色婷婷啪啪六月天| 婷婷五月天色| 老师的粉嫩小又紧水又多A片视频| 播播网色播播| 在线看的免费网站| 亚洲日本三级片| 久9精品| 丁香六月啪啪| 色情五月综合婷婷| 久久av电影| 色久五月天| 超碰在线9| 亚洲精品国产精品乱码视99| 综合一区二区三区| 五月丁香狠狠爱| 97色精品视频 | 超碰在线人妻| 婷婷久久综合| 色婷婷呢狠禁久禁| 玖玖激情网| 毛片新网地| 亚洲综合色网站| 日韩精品无码99| 精品人妻在线| 亚洲精品V天堂中文字幕| 思思久久久婷婷| 99久精品视频| 成人噜噜网| 五月天婷婷深深爱| 另类综合色| 婷婷五月天激情网| 超碰99热在线观看| 99er热精品视频| 国产xxxxx在线观看| 五月天精品综合| 色情综合| 婷婷五月天成人视频| 久久婷婷综合拍| 天天弄天天爽| 99视频在线观看网址| 99久在线精品99re8热| 热久91| 激情婷婷五月久久| 综合网五月天123| 天天五月情| 久久九九爽| 九月丁香八月婷婷加勒比| 狠狠色九月| 精品久久99| 色久五月天| 婷婷丁香五另类网站| 五月天婷婷丁香成人网| 激情久久五月天| 2025年最新亚洲在线欧美| 99性感视频| 大香蕉综合在线| 99热这里都是精品| 秋霞AV美国| 思思99re这里只有| 五月天色婷好好| www.狠狠操| 亚洲亚洲人成综合网络| enecarbon-materials.com污K127封锁请涟系@wip1688 | 国产成人精品亚洲线观看| 久久婷婷五月综合色天| 免費亭亭成人| 色碰碰视频| 色婷婷狠狠禁久久| 99久扒热| 婷婷五月天视频小说| 日本不卡中文字幕| 成人午夜在线视频| 人人干AV| 9久久婷婷国产综合精品性色| 热久综合| 殴美激情综合网| 五月婷在线播放| 五月激情影院| 九九久久腿| 99热这里只有99| www.五月瑟| 少妇高潮A片无套内谢麻豆传| 欧美婷婷日本| 精品久久99码| 国产九月婷婷| 六月婷婷狠狠做| 色综合色| 玖玖婷婷五月天毛片| 激情五月天开心网丁香无码| 色香欲综合| 欧美三日本三级少妇三99| 亚州操人在线视频| 无码成人AAAAA毛片AI换脸| 操久久精| 婷婷五月丁香网| 裸体做A爰片毛片A片免费| 无码se| 国产特黄色精品一区二区三区精品无广告 | 99A片| 国产精品久久久99视频| 久久9久| 日本91在线播放| 91人人操人人爱| 亚洲综合视频天天精品| 91成人性爱视频| AⅤ在线播放网| 六月丁香好婷婷| www.婷婷六月天| 久久九九玖玖| 激情五月影院| 五月开心啪啪| 啪啪五月天啪啪| 午夜少妇在线观看视频| 亚洲99激情| 狠狠干婷婷| 婷婷色在线观看| 五月丁香啪啪激情| 日韩黄在免| 五月丁香六月婷精品视频| 亚洲精品久久久久AV无码| 婷激情五月天视频导航| 97精品综合久久| 日日天天天| 五月丁香婷婷综合| 五月色丁香成人| www夜夜操comwww| 久久五月丁香综合| 1010日日无码| 九九99偷拍视频| 99热无码首页| 97香蕉人人在线观看| 97色色色色色| 影音先锋 一区| 激情综合网婷婷久久| 久久小说| 六月米奇色综合| 午夜不卡久久精品无码免费| 婷婷五月天日日日干干干| 色一情一乱一伦一区二区三区| 色五月婷婷五月久久| 色婷婷丁香花五月天| 五月婷六月天| 久久综合激情| 电影蜘蛛女| 五月Huangsewang| 久久aaaa片一区二区| 色婷婷aV四虎| 超碰在线综合| 久久五月天视频| 97久久五月丁香婷婷| 人妻九九九九| 久热这里只有精品性色AV| 另类视频在线| 五月丁香六月婷婷在线播放| 99视频啪啪| 噜噜噜久久| 91色色色| 国色天香成人网| 永久思思热在线| 成人无码精品1区2区3区免费看| 婷婷丁香五月精品| 激情婷婷五月社区| 久久92| 色5月婷婷色| 日本99久久| www.xtbsty.cn.com蜜乳AV| 六月激情网| 色5月婷婷| 婷婷狠狠18禁久久| 伊人五月天男人的天堂在线| 亚洲宗合激情| 激情五月天网| 久思思久视频| 婷婷五月花.97| 五月丁香啪| 亭亭五月丁香综合欧美| 五月婷婷激情| 丁香五月婷婷色| 丁香玖玖视频大全| 五月婷婷丁香五月| 亭亭五月天成人| 日日影院 | 五月天另类小说| 亚洲激情av| 精品少妇人妻AV无码专区偷人| 丁香六月欧美| 99惹在线精品免费观看| 久草免费福利视频| 激情综合九月| 五六月丁香激情视频| 色婷婷五月影视| 五月婷婷av| 丁香五月天视频在线播放| 久久精品91视频| 国产成人精品亚洲线观看| 亚洲色视频| 亚洲中文字幕网| 人碰人人人玩91| 人人草人人舔| 综合亚洲五月天| 开心久久网婷婷| 亚洲热热视频| 人人爱人人摸人人澡| 国产成人网址| 潮汕成人AV片在线| 六月丁香五月天| 激情五月天伊人影院| 性爱AV天堂| 国产操逼网站| 国产精品久久久久久久久久| 俺去也综合| wwW天天干| 色99婷婷五月天| 国产真实乱了老女人视频| 97色女人在线| 天天爽天天透天天爱| 日本三级黄色大片| 九九热99精品| 丝袜人妻| 热思思| www.色五月| 久久这里只有精品视频15| 久久久精品人妻录| 伊人五月天日日夜夜久久久天天| 中文超碰视在线| 大香蕉久久久| 丁香成人色情五月天| 国产VA播放| 色级婷婷| 国产乱子轮XXX农村| 天天狠狠色| 九九热色视频| 黄色aa观看aaguochan| 女人天堂AV| 日逼免费视频| 熟女啪啪视频| 成人五月天。COM| 色久综合| 五月天社区婷婷丁香社区| 国产亚洲精品久久久久苍井松| 99久热| 五月婷婷在线观看| 激情综合网婷婷久久| 超碰精品国产首页| 国产乱子轮XXX农村| 婷婷五月天视频免费在线观看| 一区二区中文字幕| 色欲丁香| 激情婷婷黄色五月| 色色色色色色综合网| 九月av| 婷婷五月花丁香| 丁香婷五月| 综合在线网| 91爱啪啪| 亚州综合色| 丁香色色五月| 婷婷香五月天| 九九热视频在线观看| 婷婷综合成人五月天| g00d人体西西| www.玖玖婷婷在线| 丰满少妇乱A片无码| 99色精品| 99视频自拍| 涩涩涩五月天| 颜射 精品性爱av| 综合激情肏逼网| 91AV视频| 欧美三级大片AA在线看| 99色热| 97色色视频| 久久这里都是精品| 无码激情| 人人操91| 五月婷婷亚洲色视频| 五月婷久久在线| 99热这里精| 99色在线视频观看| 婷婷网影院| 国产超碰人人| 大香av| 婷婷中文字幕网| www.刺激色网站www.| av在线观看网址| 涩玖玖免费视频| 极品少妇高潮啪啪AV无码| 日本色道视频网站| 九九综合精品| 五月香婷婷| 91chinese在线| 午夜天堂啪啪| 激情中文在线| 色五月婷婷在线| 九九99精品视频| av网址在线| 久久婷五月| 91热视频色网站| 99精品大片| 天天天天天天操| 亚洲人妻一区二区| 丁香五月激情性色郤| 97好吊操| 天天激情站| 狠狠狠色激情综合适合| 午夜成人综合| 久久婷婷五月综合啪| 精品A√| 色停停五月,在线观看| 国产精品色色| 99精品久久久久| 五月丁香婷婷基地| 中文av网| 99这里有精品视频| 亚洲五月丁香六月婷婷| 亚洲黄色精品| 久久六月天| 99精品成人无码A片观看金桔| 国产精品第一国产精品| 亚洲精品第一国产综合亚AV| 九九热99在线视频| 久久99三级在线视频| 99热网址| 超碰碰碰碰| 天天摸,天天爽| 久久婷出差欧美色两性综合网| 开心五月综合激情综合五月| 丁香六月激情综合啪啪| 激情五月天网| 五月丁香婷婷啪啪| 国产毛片欧美毛片久久久| www五月婷婷| 成人免费视频一区| 91色五月| 97碰久久| 色爱99| 中文不卡一二区| 免费精品99| 色爱亚洲| 老妇六区| 一起草av在线观看| 这里只有精品视频222| 第六色在线| 色五月六月| 五月丁香久久网| 99久热| 中文字幕人妻AV| 久久精品五月天| 26uuu.| 天天插AV丝袜中| 婷婷99中文字幕| 国产AV一区二区三区日韩| 91精品久| 久久亚洲婷婷| 天天干、天天日日| 人人干99| 北京熟妇搡BBBB搡BBBB| 韩日AV片| 欧美成人AAA片一区国产精品| 玖玖婷婷免费| 一区二区乱码视频| 97色五月天| 色五婷婷在线视频| 91狠狠综合久久| 五月色亭丁香| 91丨九色丨大屁股| 五月婷久久在线| 激情久久肏屄视频| 成人美女网| 国产亚洲99久久| 欧美综合婷婷欧美综| 另类激情首页| 99精品在线| 操人精品| 日日干综合| 激情第四色| 日本欧美成人片AAAA| 亚洲色图日韩网址| 午夜理论片最新午夜理论剧| 激情五月婷婷网| 97操操| 丁香婷婷午夜| 情婷婷五月天在线| 婷婷大乡焦噜噜| 99视频这里有精品| 99自拍网| 天天日日天天| 色五月 五月婷婷| 天天日天天爽| 大香蕉75线| 五月丁香网站在线播放| 伊人久久大香线蕉av一区| 亚洲精品又粗又大又爽A片 | 国产免费一区二区三州老师F1F1| 六月丁香婷婷开心综合基地| 影音先锋色婷婷| 97婷婷五月丁香| 婷婷色五月偷拍| 五月色色网| txt五月激情四射网综合俺也来了| www.sd-xiangsu.cpm| 久久资源网五月婷| www.色情五月天.com| 天天色色婷婷| 五月天婷婷色| 日韩久久色| 色久丁香五| 国产精品电| 99在线观看视频精品| 九九色图| 精品久久婷婷| 激情五月天伊人av| 91九色最新视频| 玖月婷婷爱丁香| 亚洲综合久| 99玖玖人人| 99惹精品视频| 成人精品一区日本无码网 | 性无码专区无码| 婷婷激情五月| www.黄色片-久久成人国产精品在线播放-999AV | 五月天大香蕉视频| 亚洲五月天婷婷综合| 五月婷婷丁香六月| 婷婷五月天最新综合你懂的 | 99视频在线| 99在线观看精品| 亚洲AV无码成人精品区电影网| 五月丁香无码| 婷婷五月在线影院| 激情亚洲五月| 99噜噜噜在线播放| 开心久久五月天| 狠狠操综合| 人人操人人操919999| 新激情五月天天在线网| 婷婷伊人五月丁香天堂网| www.ppypp| 一本色道久久综合狠狠躁一二三| 婷婷黄色五月天在线视频| 日日噜噜夜夜狠狠久久丁香六月| 六月丁香大香蕉| 亚洲无码成人| AV动漫不卡无码免费| 99爱在线视频| 婷婷五月丁香香蕉| 亚洲九九视频| 欧美毛卡| 99精品视频网| 婷婷五月激情五月丁香五月| 色欲久久久久| www.com在线操视频免费观看| 久久99美女精彩视频| 伊大人久久| 99热这里只有精品最新| 91久久久久| 99免费视频网| 激情五月瑟瑟| 亚洲色五月婷婷| 婷婷娌伦网| 无码橾| 日韩在线视频中文字幕| 亚洲不卡| 综合五月婷婷| 天天综合网在线| 人人爱干人人爱草| 色色热| 激情丁香网| 好吊兆人妻| 丁香 婷婷 亚洲 熟女| av人人操| 99久视频| 人妻AV在线观看| 97色碰| 激情综合4月| 亚洲欧洲中文日韩久久AV乱码| 丁香久久五月天视频在线观看| 日韩精品视频中文字幕| 丁香婷婷中文字幕| 色婷婷亚洲在线观看| av在线免费网站| 久9久9久9久9久9久9| 色色色色色色色色综合网| 婷婷五月天美女| 久草热久草在线视频| WWW.开心五月天.COM| 国产高潮白浆一区二区| 99热国产免费| 婷婷九九| 五月丁香六月婷| 99精品在| 亚洲无码猫咪| 丁香九月综合激情| 黄色99视频| 精品人妻午夜一区二区三区四区| 26uuuuuuuu国产| 久久九精品| 婷婷五月天伊人网| 五月香蕉婷婷| 欧美人人草草| 超碰在线综合| 99色这里| 色婷婷丁香花五月天| 综合激情站| 黄网免费看| 777精品久无码人妻蜜桃| 国精产品一区二区三区| 中日韩狠狠色| 99热这里只| 激情五月丁香在线观看直播| 激情婷婷。| 丁香六月啪啪啪| 99热思思在线观看| 99在线看视频| 超碰人人超碰| 五月深爱婷婷| 欧美碰碰| 99热免| 五月丁香欧美| 操人久久| 噜噜噜狠狠色综合| 婷婷丁香人妻天天久久| 99热亚洲| 日本欧美成人片AAAA| 婷婷六月激情| 人人草成人视频| 1024人妻无码中文字幕| 激情六月丁香| 丁香五月婷婷激情123| 9超碰在线| 99久久久久| 五月丁香六月婷婷综合网站 | 激情色五月天| 日韩久久色| yazhouzonghesese| 亚洲乱码w在线观看| 亚洲va综合va国产va中文| 五月婷婷天| 久久五月天激情美女| 任你爽在线视频| 狠狠色婷婷777| 婷婷丁香激情五月天色色| www.精品99| 狠狠噪| 色噜噜狠狠色综合日日| 成人性爱精品视频| 丁香婷婷性爱| 国产精品男人AV不卡| 久久98| 九一99| 天天干,天天舔| 五月婷婷,六月婷婷| 久色网| 五月丁香成人小说| www.五月激情.com| 丁香五月婷婷Av| 免费无码毛片一区二区A片| 色欲五月婷婷| 久久人人人人妻| 九九色图| 极品嫩草| 色色色五月婷婷| 欧美激情综合色综合啪啪五月| 五月婷免费视频久久久| 亚洲天堂大香蕉| 99九九这里有免费视频| 久久视频这里有精品99| 99这里只有精品视频| 色婷婷国色天香综合| 97干网站| 超碰人人妻| 天天日天天插| 色婷婷成人网| 五月丁香婷婷成人综合网| 伊人五月婷婷| 久久久久8888| 婷婷五月综合色小姐小说| 亚洲精品亚洲人成人网| 久久久天天啊| 丁香五月亚洲婷婷| 99热这里只有精品26| 一级二级色大片| 久久九九热re6这里有精品| 丁香六月婷婷综合| 91青娱乐青青草| 天天综合网网欲色| 色婷婷丁香五月高清在线| 91人妻人人做人碰人人爽九色| 五月婷婷 自拍| 潮汕成人AV片在线| 97韩国久久电影院| 在线视频色五月| 激情五月天啪啪| 九月婷婷在线观看| 免费观看欧美成人AA片爱我多深 | 日韩一66精品| 色狠狠综合| 97婷婷在线| 精品国婬伦V无码久久久| 中文幕无线码中文字蜜桃| 啪啪啪综合网| 类似婷婷激情综合网站| 这里只有精品视频一区| 五月婷久久| www.日本久久videos| 色女人久久| 色九月综合网| 182.t午在线观看| 国产成人精品一区二三区熟女在线 | 开心丁五月| 色色丁香| 五月丁香无码| 婷丁香五月天| 色色色色色色色色色色色色色五月天| 婷婷综合中文| 色婷婷成人做爰A片免费看网站| 九九热精品视频| 九九九九毛片| 久久性操| 激情99| 色色热| 久久怡红院| 久久久精品人妻| 激情五月天激情小说| 99综合激情久久精品久久| 五月婷婷婷婷婷婷艺术| 天干干夜夜操| 激情5月婷婷狠狠干| 五月天色婷伊人| 五月激情婷婷图片基地| 婷婷色综合| 97天堂| 99热只有精品综合| 亚洲综合字幕色色| 天天艹天天综合网| 六月丁香婷婷五月天| 97人人草| 九九国产精视频| 久热这里只有精品3| 这里只有精品偷拍| 99re熱| 色情·com| 五月婷高清视频| 五月婷婷天天色| 五月丁香综合精品| 丁香婷婷久久老熟女综合网| 人人舔天天|