免费国产精品自在自线-91精品国产色综合久久久浪潮-99热久久免费频精品-国产精品国模在线观看-久久亚洲国产精品成人?V秋霞-久久国产一级A片免费播放-亚洲国产欧洲综合97久久-久久国产白嫩美女呻吟高潮

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
返回資訊列表
天天综合五月天| 99视频内射三四| 成人精品视频99在线观看免费| 大香蕉啪啪啪| 97人人干| 这里只有精品69| 影音先锋 萱萱| 色丁香五月综合网| CAoub青青超碰| 婷婷激情六月| 噜噜噜噜婷婷五月天| 一起草性爱不卡视频| 99热老司机| www.99热视频| 开心五月深爱激情| 午夜九九九九九九九九九九九九九| 六月色婷婷| 丁香婷婷五月综合影院| www.婷婷五月天| 久久人人超| 色五月五月婷婷| 91啪啪啪啪| 色欲影香| 五月成人丁香av91| 天天爽天天操| 久久婷婷五月草视频在线播放| 少妇人妻偷人精品无码视频新浪| 九九热只有精品| 久久久27操| 色~性~乱~伦~噜| 婷婷五月深情丁香深爱日韩| 最近中文字幕大全免费版在线| 色婷婷伦理| 久久精彩视频18| 天天干 夜夜爽| 九热久| 亚洲综合色婷婷| 五月丁香手机在线| 五月婷天堂视频| 99ri精品在线| 久久这里在精品视频| 日本激情综合| 欧美 日韩 成人在线| WWW.夜夜操.com| 综合久久十三| 狠狠操狠狠| 婷婷五月图片小说网| 五月婷婷婷综合网| 99久操视频| 丁香五月婷婷视频| 色五月婷婷在线视频| 182tv992tv人之初午夜免费观看| 五月天成人小说| 91婷婷五月天综合视频| 六月丁香狠狠爱| 丁香色婷婷| 99精彩视频在线观看| 看逼中文字幕| www.色婷婷| WWW,婷婷,COM| 亚洲色婷婷激情| 六月婷婷AV| 久久婷婷五月天| 激情www| 日韩精品AV一区二区三区| 丁香婷婷免费| 九热视频在线精品15| 亚洲九N| 久久婷婷五月综合色奶水99啪| 深爱激情四射| 99re久久| 在线观看亚洲AV| 狠狠综合| 五月婷婷av| 婷婷综合五月色播| 午夜 外网 精品 在线| 99re热在线视频观看| 亚洲精品国产熟女久久久| 五月天激情视频| 亚洲瑟瑟精品在线| 九九视频这里是精品五月| 成人短视频免费观看| 成人视频网| 91九色中文字幕女在线观看| 成人丁香五月天| 亚洲激情免费久久| 久久九九怡红院| 久久视频婷婷| 综合网色综合| av五月丁香婷婷网| 夜精品无码A片一区二区蜜桃| 婷婷久久精品| jiZZdr| 天天噪夜夜爽| 天天色天天爱天天舔| 97色色婷婷五月天| 一本久道综合99| 五月丁香亭亭A片| 婷婷中文字幕| 97操碰| 婷婷伊人綜合中文字幕小说| 色婷丁香| 国产视频婷婷| 五月丁香中文| 五月天 婷 欧美亚洲| 色五月亚洲开心网| 综合五月天婷婷色| 99热这里只有精品23| 狠狠色色| 丁香六月婷婷缴情欧美| 99成人网站| 91久久久久久久| 天堂爱啪啪| 97超碰99热99| 六月婷婷青青青视频| 色色色激情网| 禁片二区| 影音先锋91资源站| 一起草无码| 91日韩美女被插视频| 97热超碰| 婷婷va| 超碰在线播放免费观看| 色五月天丁香婷婷色| 九九九AAA热视频| 97碰人人操| 五月天伊人综合| 五月婷婷与六月丁香图片激情| 久久加勒比| 色99无码| 97超级操操| 4438亚洲欧美| 性爱七区| 欧美三级欧美一级| 国产99精品免费视频| 日韩无码一区二区三区四区| 激情丰满熟妇五月| 日韩精品二三区| 婷婷夜夜操| 色情婷婷五月天| 婷婷成年人免费视频| 蜜桃五月天| 丁香五月婷久久| 996日日爱| 激情五月丁香五月| 人人操日| 中文字幕丰满孑伦无码专区 | 91ncm视频| 99热精品在线观看| 日韩aaaaa| 亚洲天堂久久| 深爱激情综合| 激情开心五月天| 99思思热只有在这里看| 九九色院| 99热大全在线观看| 五月天激情中文字幕| 五月丁香六月婷婷久久久综合| 五月婷婷丁香色吧网| 久久九九99字幕| 亚洲 在线 性爱| 97好吊操| 色人久久| 亚洲三A| 9热在线| 嫩草视频。| 亚洲热久| 五月天婷婷午夜丁香| 五月花成人网| 五月丁香五月婷婷| 亚洲激情四射| 色五月激情综合网站| 精品久久婷婷| 婷婷一本和五月丁香| 五月情婷婷| 国产免费一区二区三区三州老师F1F1.CC | 久久综合五月婷婷| 日韩欧美成人片| 婷婷爱在线观看| 婷婷亚洲在线| 亚洲黄色影视| 色约约视频一区二区三区四区五区 | 色欧洲| 丁香六月激情综合| 五月婷婷伊| 玖玖婷婷色| 婷婷色亚洲| 热99这里只有精品视频| 思思热久久艹| 久久大国产香蕉| 五月婷久久草| 呦呦视频无码播放| 色色色国产| 人妻熟人中文字幕一区二区| 六月婷婷av| 怡红院院久久| 五月婷婷综合在线| 国产黄色在线| 五月婷丁香| 六月色国内综合| 婷婷玖玖五月天| 97人人干| 97色色综合| 五月婷婷丁香五月亚洲色| 中文字幕婷婷五月天在线观看| 欧美婷婷综合| 97精品人人A片免费看| 九九热99久久99| 色,激情五月天| www.色五月.com| 综合网亚洲| 91啦丨九色丨刺激中文| 五月婷婷成人网首页| 五月天国产| 五月激情综合网| 久久思思99| 五月青青草综合| 狠狠色综合网站久久久久| 熟女激情网| 99日热在线视频| 99久久精彩视频。| 男人天堂99| 日韩在线视频9色| 五月天激情黄色小说在线观看| 玖玖资源在线视频| 五月四色激情| 丁香五月婷婷综合视频| 婷婷五月天丁香花| 操日视频| 久久久天天啊| 精品导航在线x不卡| 在线五月色播| 久久伊人五月天| 五月 激情视频| 天天操天天国产三级片处女学生妹| 華人性愛AV在線| 婷婷五月综合社区| 激情文学第四色婷婷丁香五月| 丁香婷婷成人在线播放| 国产在线激情视频| 五月婷婷干| 99在线精品视频| 4399在线日本A片| 婷婷午夜天| 免费黄网不卡AV| 婷婷色婷婷亚洲成人| 天天色天天| 亚洲亚洲人成综合网络| 亚洲综合网激情小说| 久久五月综合| 337p大胆噜噜噜噜噜91Av| 97性视频| 成人丁香五月| 久久综合九九| 婷婷成人综合| 亚洲AV免费在线| 极品色丁香| 五月色网| 婷婷五月激情四月综合 | 日韩成人精品中文字幕| 99丁香五月婷| 开心五月激情网| 亚洲成人网无码| 第四色五月婷婷| 玖玖资源站中文| 国产毛片精品一区二区色欲黄A片| 91超级碰在线| 五月香婷婷| 成年视频免费观看| 182TV大香蕉| 99久久精彩视频| 五月丁香激情综合久久| 五月婷婷久久大香蕉| 五月天激情综合网俺也去| 九九综合九九| 五月激情日本在线| 噼里啪啦完整版中文在线观看 | 欧亚色色| 免费视频WWW在线观看网站| 狠狠色婷婷7777久| 99色综合| 性色99| yjzz亚洲国产| 五月天四色房丁香亭亭| 99日逼视频| 91人人操人人爱| 久久久99视频| 久久这里只有精品99| WWW,五月天| 五月丁香色婷婷| 99热九九这里只有精品10| 亚洲精品久久久久AV无码| 99在线观看视频蜜臀| 极品少妇高潮啪啪AV无码| 黄色99网| www综合久久| 欧洲一区二区| 日本精品。999| 91丁香综合| 日韩狠狠色婷婷| 久久久国产精品黄毛片| 婷婷九月丁香| 久久五月天色婷婷| 99操九九网| 五月精品免费XXX| 中文字幕性爱视频| 亚洲色婷婷五月天| 五月婷婷之美女图片| 99热九九在线| 亚洲综合九九| 婷婷色五月色| 婷婷免费精品视频| 色婷婷啪啪| 日本三级日本三级99| 亚洲久久视频| 丁香五月天堂网AV| 91色婷婷综合久久中文字幕二区| 狠狠色婷婷丁香六月| 婷婷字幕在线| 四色99久久| 久久久精久人妻| 婷婷色中文| 激情婷婷五月| 天天狠狠干| 97人人操人人| 级人人91| www.日日夜夜.com| 五月丁花六月丁香综合| 五月婷婷激情四季| 亚洲四色五月| 大地9中文在线观看免费高清| 亚洲精品午夜国产va久久成人| 999婷婷综合| 日产精品一线二线三线芒果| 色三级色三级| 狠狠色精品综合| 97精品欧美91久久久久久久| 婷婷成人av| 91在线日| 天天做天天爱天天要| 色五月琪琪| 亚洲第一成人无码A片| 激情五月婷婷综合视频| 五月丁香激情四射综合| 狠狠 久久| 97超碰人人操| 久久婷婷成人综合色怡春院| 丁香五月老师| www激情| 久热A片| 婷婷天天日婷婷| 亚洲中文乱字字幕线在永久| 日本久久99| 中国操逼99| 99热成人精品网站| 天天综合网亚洲综合网| 九九sese| 日日操夜夜操无码免费| 日韩AC在线免费观看| www.丁香五月| 九月丁香五月婷婷| 亚洲综合干| 激情综合国产| 亚洲正能量欧美| 深爱综合网| 人人操AV| 色婷婷视频综合| 天天综合亚洲综合网天天αⅴ| 亚洲av无码精品色午夜| 色色六月| 人人艹艹艹| 99九九综合久久九九| 丁香 婷婷 亚洲 熟女| mmm1717.6dbm人人爱人人操| 婷婷五月电影| 激情四射婷婷色色色| 九月激情婷婷丁香| 久热在线中文字幕色999舞| 91丨九色丨国产打屁股| 99热99久久| 丁香成人色情五月天| 91热er| 人妻中文在线| 色色婷婷五月天| 97色在线观看视频| 99热www| 噜噜视频| 热久视频| 久久这里只| 六月婷婷影院| 人人射av| 狠狠操.COM| 久久久婷丁香五月天激情综合| 久久99日本精品视频免费观看| 伊人五月婷| 精品成人无码A片观看香草视频| 色色色99| enecarbon-materials.comWu染请涟系Bao护@wip1688| 欧洲综合色| yazhouzonghesese| 九九操综合网| 人人人操B超碰| 五月丁香婷婷爱激情综合网| 天天日天天舔天天摸| 丁香五月天黄色片| 99精品久久久久久久婷婷| 久超超碰| 永久热91| 国产无人区大片| 日亚二欧美| 麻豆精品| 国产免费AV网站| 五月婷婷日本| 激情五月综合网最新| 欧美精产国品一二三区| 久久婷婷五月天| 丁香九月综合| 大香蕉九操| www99精品| 九九aV| 五月婷精品| 伊人久久婷婷| 9999热精品在线免费播放| 色色色99韩| 激情5月婷婷狠狠干| 久久精品小视频| 天天婷婷综合| 26uu| 这里只有精彩视频| 大香蕉综合| 丁香五月天啪啪| 婷婷综合五月| 激情五月婷婷色综合| 亚洲成人无码专区| 美腿丝袜AV天堂网| 伊人久久婷婷| 色五月婷婷开心| 人人色AV| 色五月婷婷亚洲最大| 九九re精品视频在线观看| 天天日日夜夜| www.yw尤物| 夜色五月天| 久久综合网桃花| 婷婷丁香五月麻豆| 好好日激情五月天| 日韩视频99| 狼人婷婷久久| 国产肥白大熟妇BBBB视频| 婷婷综合伊人| 人妻久久久久| 久色五月| 97人人干| 婷婷日韩| 日日爱678| 99色视频在线观看| 六月婷婷综合| 国产精产国品一二三在观看| 婷婷五月激情小说| 爱久综合| 五月婷婷亞洲中文| 日本久久99| AA片在线观看视频在线播放| www.五月天色色色| 五月www| 26UUU在线观看| 人人色婷婷| 天天射影院| 色婷婷精品视频在线播放| 久久天堂精品| 五月天成人网婷婷| 丁香婷婷色| 精品一二三区久久AAA片| 51精品国内探花| 激情综合婷婷| 337p大胆噜噜噜噜噜91Av| 婷婷五月天性| 色色色婷婷五月| 四川BBB搡BBB搡多人乱亂| 色噜久| 好吊操这里只有精品| 婷婷色五月在线视频| 婷婷狠狠操| 熟女强人妻一区二区三区四区无| 深夜男女福利刺激影院一区完整| 无码色色| 少妇AB又爽又紧无码网站| av无码电影| α久久| 婷婷综合五月天| 伊人狼人干| 噜噜噜噜婷婷五月天| 五月丁香六月综合激情网| WWW.99视频| 秋霞少妇AV网站| 国产精品99久久久久久猫咪| 久狠日av| 免费看欧美成人A片无码| 日本久久视频| 成人五月天视频| 激情第四色| www色五月| 六月丁香五月激情婷婷| 激情丁香五月| 日韩综合网络男女香蕉a片| 精品国产va久久久| 成人国产欧美大片一区| 五月天狠狠网站| 五月婷婷开心网| 五月激香蕉网| 亚洲操精品| 亚洲五月天综合| 欧美性爱中文字幕| 97碰碰免费.视频| 久久性视频| 99无码精品| 婷婷五月天社区| 午夜成人天堂久久无码日韩久久| 五月婷婷丁香狠狠撸久久| 人妻AV中文系列| 久久综合影院| 五月婷婷亚洲色视频| 色伊人91在线视频| 婷婷成人丁香色情基地30| 综合亚洲五月天| 91久久综合亚洲鲁鲁五月天| 91传媒无码人妻精| 五月婷婷亞洲中文| 91色综合网站在线| 色综合久久伊伊婷婷五月| 人人爱人人草| 色丁香五月天射婷婷爱婷婷| 五月丁香| 国产成人网址| 欧美25p| 六月丁香av| 婷婷五月色色| 伊人www22综合色| 天天在线XXX| 久久日韩婷婷五月| 激情五月综合| 欧美日本黄色| 色色色综合网| 怡红院成人AV| 婷婷五月色惰| 午夜爱插插| 超碰资源在线| 亚洲av综合网| 婷婷成人视频| 电影蜘蛛女| 国外亚洲成AV人片在线观看| 丁香六月激情综合网| 丁香五月天激情综合| 97色色色色| 玖玖婷婷五月天| 欧美日韩AAAAA| 疯狂做受XXXX高潮A片| 婷婷丁香五月天色播网站| 国产激情久久久| 丁香五月天无码AV| 深爱五月天 开心网| 不卡的AV网站| 国产欧洲欧洲精品久久| 97欧美在线| 欧美丁香五月97色| 99热这里有精品首页10| 九九综合精品| 久久精品爱爱| 69精品人人人人人人人人人| 精品草原久久视频| 国产91九色| 97色久| 色综合久久99色| 五月天天天色| 五月丁香在线国产| 97人人操人人操人人操人人| 五月天激情久久| 久久综合五月天| 激情五月天网站| 丁香八月综合激情| 国产偷人爽久久久久久老妇APP| 综合色色婷婷| 成人永久免费视频在线观看| 六月婷婷综合| 香蕉AV777XXX色综合一区| 五月婷婷综合热| 一本色道久久88加勒比| 五月天激情久久| 综合色99| 天天干人人奸97| 天天摸天天舔天天天天爽| www.99久久久| 变态 另类 在线| 色五月人妻| 2015好吊操| 婷婷香蕉精品| 99玖玖免费视频| 性小说五月天| 久艹伊| 996er热| 色六月视频| 伊人婷婷大香蕉| 国产欧美日韩综合精品一区二区| 丁香五月天五码婷婷| 色色丁香五月| 大地9中文在线观看免费高清| 色婷婷五月网| 91综合国免费久入| 欧美大片| 激情av在线| 日日操夜夜爽天天天| 婷婷色五月综合| 欧美色色色色色色| 欧美成人猛片AAAAAAA| 天天色天天| 天天色播| 97操在线资源| 八戒青柠影视剧在线观看| www激情| 91艹人| 日韩淑女人妻luan伦激情精品一区二 | 五月天婷婷色| 婷婷五月五月丁香| 色激情五月| 99热首页在线30| 夜夜资源站| 99欧美热| www色色色com| 51国精产品自偷自偷综合| 久色视频首页| 久久综合网桃花| 怡红院AV亚洲一区二区三区H | 99视频久久| av不卡网站| 婷婷六月综合激情| 精品爆操| 亚洲va欧洲va国产va不卡| 日本本土色网第一区| 伊人久久中文网| 99久久人妻精品无码二区| 91精品婷婷国产综合久久| 白人荫道BBWBBB大荫道| 天天精品视频免费观看| 六月丁香综合999| 91狠狠色| 色色丁香五月| 欧美情色电影一区二区| 色五月首页| 成人综合伍月天| 色婷婷国色天香综合| 丁香五月首页| 九九色大香蕉| 97综合视频在线| 伊人综合网站| 婷婷五月激情丁香| 婷婷五月天堂| 综合久久五月天| 亚洲九九在线| 五月丁香六月婷婷综合网| www.97碰碰com| 99久久高清视频| 亚洲婷婷性爱| 免费无码毛片一区二区A片 | 狠狠做五月婷婷| 婷婷丁香色情| 色婷婷丁香| 丁香五月冃欧美| 五月婷婷色丁香| 六月丁香成人| 色五月成人婷婷| 色色婷婷五月| 香蕉视频91| 五月婷婷啪啪啪啪| 婷婷五月深深爱| 色高清无码视频| 色在线五月天免费| 婷婷五月天,影院| 亚州欧美国产久精国产99综合视频| 色综合婷婷| 97干免费视频| 九九综合| 亚洲Va成人| 性爱动图国产麻豆一区二区三区 | 少妇人妻丰满做爰XXX| 国产成人精品亚洲线观看| 色婷久久| 99热这里有精品| 五月丁香黄色视频| 五月激情偷拍| 五月天婷婷基地| 九月色婷婷| 色五月激情五月| 丰满人妻妇伦又伦精品国产| 另类A片| 中文字幕在线观看视频www| 91性高潮久久久久久久久| 九九色热| 久久五月天色婷婷| 天天干天天做| 色色激情五月天| 婷婷丁香五月天激情| 91vip在线观看| 大学生高潮无套内谢视频| 性做久久久久久久免费看| 无遮挡国产高潮视频免费观看| 精品国产乱码久久久久夜深人妻| 日韩黄在免| 9999热免费视频视频| 午夜av网| 网色99| 天天爽天天爽天天爽天天爽天天爽| 丁香香五月激情免费视频| 天天拍夜夜爽日日| 一起肏在线视频| 日韩无码成人电影| 韩国97天堂| 丁香狠狠色婷婷久久无码视频| 二区成人视频| 五月天天堂久久| 思思视频久久| 五月丁香天天| 久久成人人妻| 97精品欧美91久久久久久久| 99成人| 午夜天天精品视频| 九色在线观看91av| 噜噜噜噜噜久| 久久99久久99精品免观看软件| 色XX综合网| 婷婷五月天久久| 91精品久久久久久77777| 99热这里只有精品9| 韩国中文字幕91| 久久婷婷五| 久久99免费视屏| 久草大| 久久五月天黄色五月天色网址| 五月激情综| 九月婷婷综合| 人操综合| 九九色院| 深爱五月中文字幕| www.狠狠操| 色色色丁香| 亚洲AV网址| 狠狠搞综合色| 九九热免费视频| 9久精品| av操一操| 丁香五月色播中文在线播放| 天天艹夜夜艹| 少妇高潮呻吟A片免费看软件| 婷婷色一二三区波多野结衣| 狠狠色噜噜狠狠色噜噜噜999| 丁香五月天的网址。| 91操屁股| 操笔无码| 婷婷成人AV| 中文中文在线| 国产五月天婷婷| 天天狠狠色噜噜| 丁香五月婷婷AV| 大香蕉99| 成人国产网| 人妻性爱av网站| 五月丁花六月丁香综合| 操操碰| 99热热这里只精品996小说| 五月的婷婷六月丁香| 九色91视频| 97超级碰人人| 欧美婷婷丁香五月社区| 丁香五月天在线| 日日夜夜狠狠| 亚洲第一成人无码A片| 这里只有精彩视| 99热精品在线| 色婷婷大香蕉| 五月天婷婷色综合| 欧美性生交xXxX久久久| 欧美婷婷| 五月久久婷婷天堂视频| 激情综合五月丁香六月婷婷| 深爱五月激情网| 9久国产精品| 4438全国最大视频成人网站在线观看| 思思99精品视频在线观看| 婷婷综合色图| 国产婷婷久久| 九九热视频免费| 天天色粽合合合合合合合| 大香蕉狼人久久| av中文网站| 亚洲综合激情五月久久| 日本三级黄色大片| 五月婷婷婷丁香播| 99re视频在线精品| 六月婷婷色宗合| 99综合| 色99在线视频| 无码色| 激情五月五月五月婷婷| 天天综合网在线| 丁香五月在线自慰| 五月天色婷婷伊人网| 嫩BBB搡BBBB榛BBBB| 激情五月天婷婷| 九色91国产| 99色精品| 色情成人五月天| 色yeye欧美| 人人视频色| 狠狠爱深色婷婷综合| 九九综合精品| 六月激情网| 五月天婷婷色在线视频免费观看 | 99热亚洲精品| www.婷婷久久五月天| 亚洲 无码 中文字幕 中出| 99这里有精品视频| 精品久久人妻热| 日本色色色| 99色中文| 日本精品99网站| 91精品无码| 综合色色网| 91精品婷婷国产综合久久| 日韩成人AV在线播放| av在线中文| 久久久精品AV| 91超碰在线观看| 午夜免费试看| 9久热| 五月激情六月丁香| 亚洲视99| 免费看欧美成人A片无码| 99久久国产宗和精品1上映| 色五月婷婷亚洲| 男人的天堂999| 99色在线视频| 色五月丁香激情视频| 色135综合网| 亚洲男人的天堂婷婷色五月| 伊人五月婷| 天天操狠狠操| 97色视频网| 97色色综合| 九月影院義母在线播放| 91凹凸在线| 99九九在线精品热动漫| 9久久精品| 99热精品在线| 色婷婷色情| 五月天另类小说久久小说网| 人妻久久久久久久 | www.99在线| 一区二区三区视频| 婷婷五月丁香五月天| 九热网站| 五月天俺去也| 色播五月综合网| 丁香婷婷大香蕉| 综合激情开心五月| 超碰人人摸人人操| 国产激情婷婷| 婷婷色五月丁香六月欧美啪| 97精品综合久久| 26uuu淫色| 国内自拍97在线| 婷婷五月天激情丁香| 五月婷婷导航| 99色视频在线| 激情五月天综合图片小说网站| 久久久久9久无码视频| 九九精品免费| AV在线中文| 五月婷在线| 五月六月激情| 五月丁香六月综合基地| 色婷婷激情| 五月停停激情网| 婷婷五月综合久久中文字幕| 思思99热| 婷婷激情社区| 综合激情婷婷| 亚洲一区二区色图-亚洲精品国产精品乱码-成人AV | 亚洲综合久| 婷婷五日b| 伊人狼人干| 超碰成人在线观看| 婷婷玉月丁香五月在线视频| 97色视频网| 六月婷婷五月丁香| 婷婷五月天在线观看免费| 国产精品操| 国产午夜精品一区二区三区四区| 天天干天天日天天操| 色婷婷操逼| 《亚洲操B久久免费在线观看,亚洲操B久久在线播放》在线播放 - 高清资源 - 97 | 婷婷五月丁香色色| 激情久久肏屄视频| 色情激情五月| 五月婷婷色影院| 色五月婷婷成人| 日日干天天射| 婷婷成人综合| 大香久久伊人网| 激情综合网五月婷婷| 五月天婷婷影院影院观看| 性韩日色婷婷五月天激情啪啪XXX| 日本在线观看99| www久| 国产精品热搜丁香五月婷婷| 99,色| 久久99日本精品视频免费观看| 日本va欧美va国产激情| 嫩草乱码一区三区四区| 丁香婷婷五月天色综合| 日本一区二区三区精品视频| 日本综合久久| 五月亭亭六月天| 青青操avbb| 在线观看av网站| 九月丁香八月婷婷加勒比| 丁香五月婷婷啪啪啪| 深夜A片| 另类的婷婷| 五月丁香成年黄色| 丁香五月手机在线| 婷婷第六色| 在线伦子99热| 五月婷婷操操| 九九九九毛片| 舔色婷婷| 色情五月天se| 色色婷婷五月| 视频一区二区在线| 中文人妻AV久久人妻18| 五月天成人综合| 五月丁香伊人网| 六月婷婷av| 91九色无码内射| 成人无码免费一区二区中文| 激情影院丁香五月| 婷婷五月天激情电影| 五月婷婷网五月在线| 丁香狠狠| 丁香五月影院| 综合久久五月天| 五月婷婷之综合激情在线| 少妇出轨做爰高潮A片| 综合色色色| 成人五月天丁香| 色婷婷影音| 丁香五月天激情五月天激情五月天激情网| 丁香五月中文字幕| 五月激情射| 99资源在线| 色婷婷精| 91久女| 日韩av一区二区在线/日产精品久久久| 99爱视频免费| 色婷久九| 婷婷综合成人| 99只有这里是精品| 襙比视频| 丁香5月综合啪啪| 乱女乱妇熟女熟妇综合网站| 五月丁香久久| 五月天操逼网| 另类激情四射| 99色在线观看视频| 六月婷婷九月丁香| 久久综合色情网站| 97操男人的天堂| 色综合9| 嫩草视频。| 综合另类激情| 国产性爱色| 玖玖婷婷视频| 婷婷色啪| 婷婷丁香五月综合| 人妻自慰在线| 色五月激情五月丁香五月婷婷啪啪综合 | 色99视频| 激情五月激情综合俺也去婷婷小说| 这里只有精品偷拍| 丁香婷婷免费| 五月丁香啪啪综合网| 天天爽人人爽| 26uuu亚洲欧美| 国产精品婷婷午夜在线观看| 天天爽天天做| 色五月综合在线| 婷婷六月色| 亚洲操b| 99热官网| 免费婷婷| 综合网五月| 99爱精品| 亚洲人妻av| 福利视频在线播放| 日韩成人影片网站| 国成人网| 99热99| 国产9色在线/日韩| 丁香久久| 婷婷五月情色| 激情五月婷在线精品| 99精品小视频| 久久综合首页| 九热视频| 五月天开心网| 色色婷婷综合| 一级七香蕉| 五月婷婷深深爱| 久久99这里只有精品| 色五月激情综合| 五月丁香六月在线| 九九热只有精品| 国产在线网| 国产99久9在线| AV九九| 日日夜夜天天爽| 97色在线观看视频| 九九热免费视频| 黄色AAAAAAA| 五月激情视频| m色激情网| 综合网啪啪| 色五月情| 91精品国产日韩91久久久久久国模| 亚洲va日| 激情五月天伊人av| 色色婷| 99只有精品9| 六月婷婷av| 天天舔天天插天天干| 丁香五月综合激情性爱| av亚洲国产小电影| 五月天婷婷丁香花| 激情五月婷婷欧美极品| 丁香五月手机在线| 五月天色色色| 玖玖五月丁香| 日日色五月天| 色婷婷先锋| 丁香香蕉射射射| w婷婷五月婷婷w| 欧美97p| 久操热线| 99碰碰碰| 丁香五月在线伊人| 天天干天天干天天干天天干天天干| 思思热在线| 国产在线自| 678五月丁香亚洲综合| 日韩超碰在线| 丁香伍月婷电影全集| 成人丁香五月| 久久精品视频在这里有| 久久激情网| 黑人无码一区| 99精品无码网站| 色九月| 月色色综合婷婷网| 无码G高清天| 91操网| ww超碰在线| 亚洲五月天天| 亚洲五月婷| 色五月婷婷91在线| 欧美综合激情五月丁香| 婷婷五月天堂| 久久精品日| 婷婷五月天色网久| 丁香五月天网站| 色五月激情综合网站| 激情开心五月婷婷| 第四色激情网| 亚洲精品成人片在线播| 五月欧美色色五月| 色婷婷六月天| 91婷婷色五月| 色婷婷亚洲五月天| 91干婷婷| 狠狠五月丁香色婷| 精品久久99码| 午夜丁香综合婷婷| 婷婷丁香人妻天天| 丰满少妇乱A片无码| wWwCom夜操wwW| 五月天丁香婷| 9久久久| 色女伊人| 婷婷五月丁香亚洲| 激情六月五月婷婷综合网| 26uuu欧美日本| 人妻激情在线| 五月丁香影视| 婷婷五月激情视频网| 亚洲性色XXXXX| 禁欲电影完整版在线播放| 婷婷99狠狠| 九九热99视频| 日日操夜夜操无码免费| 色五月色五天免费视频| 五月天丁香婷婷视频网址| 77799热| 丁香六月激情综合| 操一区| 九九九热精品| 日本在线视频播放91| 丁香五月天堂网| 97婷婷五月激情六月丁香伊人| 精品,99| 色五月婷婷成人| 毛片新网地| 不卡在线视频| 青青草婷婷综合五月| 八戒青柠影视剧在线观看| 天天综合色丁香| 色五月婷婷色| 狠狠色婷婷在线| 能看的AV| 国产xxxxx在线观看| 亚洲操B| 黄色99视频| 激情综合五月天| 婷婷五月丁香五月丁香| 91一起操| 亚洲综合激| 少妇性按摩无码中文A片| 亚洲色五月天是什么| 免费看欧美成人A片无码| 日韩激情婷婷五月天| 国产做A爰片毛片A片美国| 五月色 亚洲| 91久久精品国产91性色TV| www.操逼comm| 狠狠久久婷五月综合色| 五月婷婷久草在线视频综合| 99久久婷婷国产综合亚洲| 天天日,天天干,天天操| 丁香五月天婷婷91| 婷婷五月电影| 五月婷婷久久综合| 丁香五月婷婷基地| 天天日狠狠| 精品人妻在线免费观看| 这里只有精品视频一区| 1024在线视频| 五月天婷婷免费视频| 另类A片| 久久人妻精品| 色狠狠六月| 91在线资源| 丁香成人五月天| 97操碰视频| 激情小说五月天| 婷婷激情五月综合基地| 婷婷九月久久| 九九无码| 91啪啪| 成人国产欧美大片一区| 九九综合久久| www.日本91| 久久婷婷综合五月天| 五月婷成人网| 狠狠色综合久久久久| 久色大| 天天爱天天做天天日| 麻豆WWWCOM内射软件| 亚洲天天综合| 色色色地址| 久久曰曰| 五月丁香婷婷开心| 久热91精品| 丁香六月激情综合| 天天色天天色天天色天天色天天色| 一本色道久久综合狠狠躁小说| 丁香啪啪| 九九综合影音先锋| 日日操夜夜骑| 久久人人九九| 五月婷婷综合激情| 狠狠色噜噜狠狠| 成人五月天丁香| EEUSS鲁片一区二区三区| 久久综合激情婷婷激情| 国产精品99久久久久久猫咪| 性爱视频久久| 婷婷丁香五月亚洲| caop在线视频| 亚洲成人影视在线| 玖玖婷婷色五月| 久久这里只有精品07| 99热欲| 99热精品在线| 特级片神马电影| 婷婷九九| 色婷婷网大全在线| 六月婷婷综合激情| 五月天天天开心激情网| 婷婷五月激情综合啪啪| 色99在线视频| 2025年最新亚洲在线欧美| 亚洲成人一区| 久99在线| 97自拍视频在线| 丁香五月天天高清在线| 色婷婷香蕉| 亚洲无码性爱| 色色激情五月天| 99热成人在线观看| 91欧美| 操b视频在线观看一区二区| 色情婷婷五月天| 99精色| 亚洲色爽| 九九大香蕉黄色影院| 亚洲色无码A片一区二区麻豆| 9月色婷婷| 天天天天做夜夜夜夜做| 人人操99| 亚洲AV电影美洲AV电影| 中文AV网| 狠狠狠色激情综合适合| 伊人成人宗合网| 色婷婷免费观看| 五月婷婷激情啪啪| 五月天婷婷久久| 天天爽天天| 色天天综合成人网| 久久久久妻| 欧美色碰| 五月丁香偷拍| 五月丁香婷婷六月天| 亚洲综合久| 色色五月丁香婷婷综合| 黄桃AV无码免费一区二区三区| 久99视频在线观看| 人人噜天天上| 婷婷五月天激情偷拍| 丁香五月首页| 色色色色色五月| 婷婷丁香五月天综合激情| 婷婷综合激情| 夜夜躁婷婷AV| 五月婷婷免费看| 日韩av干| 性色欲情 网站| aaa久久| 色婷大香蕉| 亚洲国产精品成人午夜| 狠狠干综合| 日韩999| 97欧美在线| 久久久久er热| www激情网| 丁香五月婷婷动漫| 久久视网36| 亚洲成人AV在线观看| 一本色道久久88加勒比—| 日韩有码一区| 亚洲成AV人片在线观看|