日韩精品一区二区三区在线视频放-无码中文字幕V?一区二区-成年片免费观看视频-国内少妇人妻丰满av-国产精品中文字幕免费观看-亚洲成人久久一区二区三区-国内少妇偷人精品视频无缓冲-一区二区国产精品日本一区二区三区在线网

ARTICLE DETAIL

資訊詳情

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

PyTorch編譯優(yōu)化實(shí)戰(zhàn):torch.compile、Triton與XLA性能調(diào)優(yōu)指南

PyTorch編譯優(yōu)化實(shí)戰(zhàn):torch.compile、Triton與XLA性能調(diào)優(yōu)指南 1. 性能瓶頸到底在哪先從一份“看起來(lái)很忙”的Profiling說(shuō)起前陣子幫朋友排查一個(gè)訓(xùn)練任務(wù)A100上GPU利用率看著有八成Loss也在降但總感覺(jué)哪里不對(duì)勁。跑了一輪Profile之后發(fā)現(xiàn)實(shí)際計(jì)算核心Kernel的吞吐遠(yuǎn)沒(méi)有跑滿(mǎn)大量時(shí)間花在了小算子的啟動(dòng)和內(nèi)存拷貝上。這種“表面熱鬧、實(shí)際虛高”的利用率其實(shí)很多搞PyTorch訓(xùn)練的人都遇到過(guò)——模型代碼不是瓶頸PyTorch本身的執(zhí)行機(jī)制才是瓶頸。這也正是我當(dāng)時(shí)決定系統(tǒng)整理這套《AI系統(tǒng)性能工程學(xué)習(xí)筆記》的原因所在而第十四篇的內(nèi)容就是圍繞PyTorch的編譯優(yōu)化這條主線拆開(kāi)講講我在實(shí)際項(xiàng)目中反復(fù)驗(yàn)證過(guò)的三個(gè)關(guān)鍵方向。要理解torch.compile的價(jià)值首先要理解PyTorch默認(rèn)的“逐算子執(zhí)行”模式Eager Mode是怎么回事。簡(jiǎn)單來(lái)說(shuō)你用模型定義寫(xiě)出一層卷積、一層歸一化、一個(gè)ReLUEager模式下PyTorch就會(huì)按順序把每個(gè)算子逐個(gè)交給GPU執(zhí)行。每次執(zhí)行都涉及一次Python層的調(diào)用、一次GPU Kernel的Launch、一次數(shù)據(jù)的搬運(yùn)。小算子和短Kernel特別多的時(shí)候啟動(dòng)開(kāi)銷(xiāo)占比就會(huì)直線上升。哪怕每秒鐘能提交幾萬(wàn)個(gè)KernelGPU真正的計(jì)算單元卻經(jīng)常處于“等任務(wù)”的狀態(tài)。那為什么不把所有算子一次性做完因?yàn)樗阕又g存在數(shù)據(jù)依賴(lài)。卷積的輸出要喂給歸一化歸一化的輸出要喂給ReLU每一步都不能跳過(guò)??扇绻茉诒WC依賴(lài)關(guān)系不被打亂的前提下把多個(gè)連續(xù)算子融合成一個(gè)大的Kernel那啟動(dòng)開(kāi)銷(xiāo)就能被壓到很低內(nèi)存讀寫(xiě)也能少幾個(gè)來(lái)回。這個(gè)思路聽(tīng)著不難但落地起來(lái)步驟非?,嵥椤纫治鲇?jì)算圖又要生成高性能Kernel還要在不同硬件上做適配。PyTorch社區(qū)的答案是torch.compile它把這條鏈路做成了一行代碼就能用起來(lái)的方案。坦白說(shuō)torch.compile剛出來(lái)的時(shí)候我是持觀望態(tài)度的。畢竟PyTorch一直以靈活和動(dòng)態(tài)圖著稱(chēng)硬上編譯優(yōu)化很容易破壞調(diào)試體驗(yàn)。但實(shí)測(cè)了幾批CV和NLP模型之后我的態(tài)度發(fā)生了明顯轉(zhuǎn)變?cè)诓恍枰l繁改圖、動(dòng)態(tài)Shape不嚴(yán)重的訓(xùn)練場(chǎng)景下torch.compile帶來(lái)的提升相當(dāng)直觀尤其是在A100、H100這類(lèi)新架構(gòu)上收益常常能到20%到50%。更關(guān)鍵的是它解決的是“全局優(yōu)化”的問(wèn)題而不是零敲碎打地優(yōu)化某個(gè)算子。從這一篇開(kāi)始我會(huì)把torch.compile的機(jī)制、Triton在其中的作用、以及XLA作為另一條編譯器路線的取舍完整地串起來(lái)講。一個(gè)必須明確的點(diǎn)是torch.compile不是銀彈。它適合那些結(jié)構(gòu)穩(wěn)定、算子類(lèi)型清楚的模型如果模型里到處都是Python控制流、動(dòng)態(tài)Shape、自定義算子編譯優(yōu)化能覆蓋的區(qū)間就會(huì)被壓縮得很厲害。所以這篇文章不只是教你怎么用也會(huì)告訴你什么情況下“不要用它”以及在XLA和torch.compile之間到底該怎么選。2. 拆開(kāi)torch.compile的引擎蓋Dynamo、Graph Break 與 Inductortorch.compile能一行代碼接管優(yōu)化本質(zhì)上是因?yàn)樗鼉?nèi)部有一條流水線式的處理路徑。理解這條路徑比死記API參數(shù)有用得多。我會(huì)從自己調(diào)試過(guò)的實(shí)際案例出發(fā)把它的工作過(guò)程拆成三個(gè)階段講。2.1 Dynamo如何在Python的動(dòng)態(tài)世界里“抓到”計(jì)算圖PyTorch模型是用Python寫(xiě)的Python太靈活了一個(gè)if條件、一個(gè)for循環(huán)、一次字典索引都有可能在運(yùn)行時(shí)改變計(jì)算圖的結(jié)構(gòu)。真要等模型跑完再拿到完整靜態(tài)圖那延遲就太大了。torch.compile在0.x到1.x的迭代中逐步用Dynamo作為前端核心思路是在“保證動(dòng)態(tài)行為正確”的前提下捕獲盡可能大的計(jì)算子圖。Dynamo的做法是追蹤字節(jié)碼。它會(huì)攔截Python函數(shù)的執(zhí)行過(guò)程記錄哪些操作是Tensor計(jì)算哪些操作是純Python邏輯。對(duì)于Tensor計(jì)算Dynamo會(huì)把它轉(zhuǎn)換為計(jì)算圖節(jié)點(diǎn)對(duì)于Python邏輯如果無(wú)法翻譯成圖節(jié)點(diǎn)就標(biāo)志為“Graph Break”。Graph Break之后模型執(zhí)行會(huì)退回Eager模式等跑到下一個(gè)可編譯區(qū)域再重新進(jìn)入優(yōu)化路徑。Graph Break不是報(bào)錯(cuò)但它的多少直接決定了優(yōu)化效果。如果一個(gè)模型里有幾十個(gè)Graph Break那torch.compile幾乎等于沒(méi)優(yōu)化因?yàn)榇蟛糠謺r(shí)間都留在了解釋執(zhí)行狀態(tài)。我在實(shí)戰(zhàn)中遇到過(guò)一個(gè)比較典型的例子模型里對(duì)某個(gè)Tensor做了.item()操作然后根據(jù)這個(gè)標(biāo)量去決定是否執(zhí)行某個(gè)分支。這種寫(xiě)法在調(diào)試時(shí)很自然但Dynamo會(huì)在這里斷掉后面一大段分支都變成Eager路徑。解決辦法也很直接——把.item()移到模型外部或者在損失函數(shù)中避免使用同步操作。你可以在每次編譯后查看torch._dynamo.explain的輸出它能夠列出Graph Break的位置和原因這一招在排查性能問(wèn)題時(shí)極其有效。2.2 Inductor從計(jì)算圖到GPU Kernel的“翻譯官”Dynamo抓到計(jì)算子圖后接下來(lái)的工作就交給后端編譯器。PyTorch默認(rèn)的生成后端是Inductor。它的職責(zé)是把計(jì)算圖翻譯成高性能的GPU Kernel代碼——在NVIDIA GPU上Inductor會(huì)把算子融合和代碼生成的任務(wù)進(jìn)一步交給Triton去搞定。Inductor采用的是基于IR的重寫(xiě)方式。它會(huì)先讀入計(jì)算圖嘗試做元素級(jí)融合、Reduce融合、點(diǎn)乘融合等操作。比如說(shuō)對(duì)一個(gè)Tensor先做x 1再乘2再用tanh激活這三步在Eager模式下至少三四個(gè)Kernel但在Inductor里會(huì)被融合成一個(gè)Triton Kernel一次讀寫(xiě)就完成全部計(jì)算。GPU的內(nèi)存帶寬往往是最大約束少一次全量讀寫(xiě)收益就非常明顯。在torch.compile的配置中backend參數(shù)控制使用哪個(gè)編譯器后端。默認(rèn)是inductor但也可以切換為cudagraphs、tvm等。我實(shí)際用下來(lái)Inductor在NVIDIA卡上的兼容性和性能綜合表現(xiàn)最好。cudagraphs的思路是把一系列Kernel的啟動(dòng)信息錄制下來(lái)然后重復(fù)回放減少CPU端的啟動(dòng)開(kāi)銷(xiāo)但對(duì)算子融合無(wú)能為力。如果你的模型本身就是大算子為主瓶頸不明顯那cudagraphs可能就夠了如果模型是小算子密集型的老老實(shí)實(shí)用Inductor。還有一點(diǎn)容易踩坑torch.compile默認(rèn)會(huì)嘗試動(dòng)態(tài)Shape的支持但開(kāi)啟動(dòng)態(tài)Shape等于放棄了一部分融合優(yōu)化。如果你的輸入尺寸在訓(xùn)練中基本固定可以考慮用dynamicFalse或者把輸入Tensor的尺寸約束住讓Inductor生成更激進(jìn)的專(zhuān)用代碼。我們做離線推理優(yōu)化時(shí)就是這么干的一個(gè)固定尺寸的模型編譯后通常能再壓掉10%左右的延遲。2.3 mode參數(shù)該怎么選default、reduce-overhead還是max-autotunetorch.compile的mode參數(shù)是一個(gè)很容易被忽略但影響很大的選項(xiàng)。官方提供了default、reduce-overhead和max-autotune三檔。default模式下Inductor會(huì)做一些低成本的優(yōu)化編譯時(shí)間短但也意味著放棄了部分更激進(jìn)的改動(dòng)reduce-overhead會(huì)在編譯后引入CUDA Graph錄制對(duì)很多小模型有額外收益max-autotune則會(huì)對(duì)生成的Triton Kernel做大量自動(dòng)調(diào)參性能上限最高但編譯時(shí)間可能長(zhǎng)達(dá)幾分鐘到幾十分鐘。我自己的建議是先用default跑通確認(rèn)無(wú)功能問(wèn)題、無(wú)Graph Break導(dǎo)致的性能回退再?lài)L試reduce-overhead。如果你的模型在多個(gè)批量尺寸上都要用不要貿(mào)然開(kāi)max-autotune因?yàn)閍utotune是針對(duì)固定Shape做的。數(shù)據(jù)增強(qiáng)或動(dòng)態(tài)批量導(dǎo)致Shape頻繁變化時(shí)max-autotune的收益會(huì)被命中率稀釋反而浪費(fèi)了編譯時(shí)間。還有一個(gè)重要技巧在A100或者H100上reduce-overhead通常會(huì)讓小批量訓(xùn)練的速度提升非常明顯因?yàn)樗鼫p少了CPU到GPU之間的同步等待。但在V100這些老卡上CUDA Graph帶來(lái)的收益相對(duì)有限因?yàn)橛布旧淼膯?dòng)延遲沒(méi)那么敏感。環(huán)境不同同樣參數(shù)跑出來(lái)的效果可能完全不一樣這也是為什么我建議任何優(yōu)化都要結(jié)合自己的硬件和模型實(shí)測(cè)而不是照搬網(wǎng)上報(bào)告的數(shù)字。3. Triton 在 torch.compile 里的角色寫(xiě)一次、到處亂跑的高性能內(nèi)核聊到Torch編譯優(yōu)化Triton是一個(gè)繞不開(kāi)的名字。很多剛接觸的人會(huì)把Triton誤解為某種第三方算子庫(kù)其實(shí)它在torch.compile中的作用更底層Inductor生成的代碼很大一部分是Triton語(yǔ)言的Kernel??梢园阉譁\地理解為“GPU上的Python”——用一套類(lèi)Python的語(yǔ)法寫(xiě)出能夠直接在CUDA設(shè)備上高效運(yùn)行的GPU Kernel而不需要手動(dòng)管理線程塊、共享內(nèi)存、同步指令那些復(fù)雜細(xì)節(jié)。3.1 Triton Kernel 到底解決了什么問(wèn)題傳統(tǒng)CUDA編程最難的不是“讓程序跑對(duì)”而是“讓程序跑快”。你要手動(dòng)決定每個(gè)線程負(fù)責(zé)哪個(gè)元素要處理內(nèi)存合并訪問(wèn)要在不同層級(jí)的內(nèi)存之間做搬運(yùn)還要處理Bank Conflict這類(lèi)隱藏很深的性能殺手。寫(xiě)出來(lái)的代碼一旦換了GPU架構(gòu)往往又要重新調(diào)整。Triton的思路是把這些底層細(xì)節(jié)抽象成塊級(jí)操作你只需要描述“每個(gè)塊負(fù)責(zé)計(jì)算什么”而塊內(nèi)部怎么劃分線程、怎么分配寄存器、怎么訪問(wèn)顯存由編譯器自動(dòng)決定。這個(gè)抽象對(duì)自動(dòng)調(diào)優(yōu)特別友好。編譯器可以在一次編譯過(guò)程中生成多個(gè)候選版本然后在真實(shí)硬件上跑一遍選擇最快的那一個(gè)。Inductor在生成Kernel時(shí)就會(huì)調(diào)用Triton做這輪調(diào)優(yōu)。你在日志里看到的“Triton kernel”字樣本質(zhì)上就是Inductor針對(duì)某個(gè)融合子圖生成的GPU代碼。實(shí)際效果上我跑過(guò)一個(gè)典型的ResNet風(fēng)格模型Eager模式下有大概600多個(gè)Kernel執(zhí)行開(kāi)啟torch.compile Inductor Triton后Kernel數(shù)量降到兩百左右端到端訓(xùn)練時(shí)間減少了35%。這個(gè)降幅不是因?yàn)槟骋粋€(gè)算子變快了而是因?yàn)榇罅啃ernel被合并成少量大KernelGPU的執(zhí)行效率因此獲得整體提升。用一句話概括就是Triton不是讓某個(gè)算子從10微秒變成1微秒而是讓幾百次啟動(dòng)從“每次都在浪費(fèi)”變成“每次都真正在計(jì)算”。3.2 Triton 在非編譯場(chǎng)景下的用法手寫(xiě)自定義Kerneltorch.compile之外Triton也可以作為獨(dú)立工具來(lái)寫(xiě)自定義算子。PyTorch里寫(xiě)高性能自定義算子傳統(tǒng)路線是寫(xiě)CUDA C擴(kuò)展然后通過(guò)torch.utils.cpp_extension編譯加載。這條路功能上限高但開(kāi)發(fā)效率低——你得同時(shí)掌握C、CUDA和PyTorch的C接口。Triton提供了一個(gè)折中方案用Python編寫(xiě).triton.kernel裝飾的Kernel函數(shù)運(yùn)行時(shí)自動(dòng)編譯成GPU代碼。調(diào)用的方式跟普通PyTorch函數(shù)一樣傳Tensor進(jìn)去就行。我舉一個(gè)實(shí)際項(xiàng)目里的例子我們要做一個(gè)自定義的注意力掩碼算子標(biāo)準(zhǔn)PyTorch實(shí)現(xiàn)里涉及多次reshape和mask操作顯存開(kāi)銷(xiāo)高。用Triton重寫(xiě)后一個(gè)Kernel內(nèi)完成了mask和softmax的部分融合顯存占用顯著下降速度也快了不少。當(dāng)時(shí)只花了半天時(shí)間就寫(xiě)完了而如果用CUDA C可能要兩三天起步。如果你打算自己動(dòng)手寫(xiě)Triton Kernel建議從簡(jiǎn)單的逐元素算子開(kāi)始練手比如把“ReLU 縮放 偏移”融合成一個(gè)自定義Kernel。把基礎(chǔ)語(yǔ)法、tl.load和tl.store的使用方式搞明白后再?lài)L試更復(fù)雜的Reduce算子。Triton的官方教程里有一個(gè)softmax例子是非常好的入門(mén)材料——同樣一個(gè)功能分別用PyTorch原生實(shí)現(xiàn)和Triton Kernel實(shí)現(xiàn)對(duì)比兩者的時(shí)間消耗你會(huì)很快理解編譯優(yōu)化的核心價(jià)值在哪里。3.3 Triton 版本兼容與安裝坑位Triton目前最常見(jiàn)的安裝方式是隨torch一同安裝。torch.compile在NVIDIA后端會(huì)依賴(lài)Triton所以你在用conda或pip安裝較新版本的PyTorch時(shí)Triton通常已經(jīng)是配套的。但如果你之前手動(dòng)裝過(guò)舊版Triton或者環(huán)境里有多個(gè)PyTorch版本很容易出現(xiàn)“torch版本和Triton版本不匹配”的警告。我在排查環(huán)境時(shí)發(fā)現(xiàn)一個(gè)規(guī)律每當(dāng)PyTorch發(fā)布新版本社區(qū)里就會(huì)出現(xiàn)一批“安裝完torch.compile報(bào)錯(cuò)找不到Triton”或“編譯Kernel報(bào)錯(cuò)版本太舊”的帖子。這時(shí)候首先要做的不是重裝Triton而是確認(rèn)當(dāng)前PyTorch要求的Triton版本范圍。最穩(wěn)妥的方案是直接使用官方推薦的安裝指令pip或conda重新裝一遍PyTorch讓依賴(lài)自動(dòng)拉齊。Google Colab和Kaggle Notebook這類(lèi)云端環(huán)境偶爾也會(huì)出現(xiàn)預(yù)裝Triton版本和torch不匹配的情況重置運(yùn)行時(shí)或升級(jí)torch就可以解決。如果你在CPU-only的機(jī)器上跑torch.compile會(huì)發(fā)現(xiàn)很多Triton相關(guān)功能不可用或者編譯速度極慢。這不是你的代碼有問(wèn)題而是Triton的GPU后端需要CUDA編譯器工具鏈。CPU環(huán)境下Inductor會(huì)嘗試生成C Kernel功能上能跑通但優(yōu)化幅度和GPU場(chǎng)景沒(méi)有可比性。所以建議在動(dòng)手實(shí)踐之前先確認(rèn)自己的環(huán)境有可用的NVIDIA GPU并且PyTorch的CUDA版本和驅(qū)動(dòng)是匹配的。4. XLA 后端另一條編譯器路線的優(yōu)勢(shì)與代價(jià)PyTorch的編譯優(yōu)化不止torch.compile一條路徑。XLAAccelerated Linear Algebra是一個(gè)從TensorFlow生態(tài)里沉淀下來(lái)的編譯器框架也可以作為PyTorch的后端使用。你只需要把模型轉(zhuǎn)換為torch_xla下的執(zhí)行設(shè)備就可以讓計(jì)算圖經(jīng)過(guò)XLA的優(yōu)化后再編譯為對(duì)應(yīng)的硬件指令。這個(gè)方案的好處是跨硬件能力很強(qiáng)——從NVIDIA GPU到Google TPUXLA都有對(duì)應(yīng)的編譯目標(biāo)。4.1 XLA 的工作原理和典型場(chǎng)景XLA的核心思路是“拿到完整的計(jì)算圖再做全局優(yōu)化”。它會(huì)把輸入的計(jì)算圖做算子融合、內(nèi)存規(guī)劃、布局優(yōu)化等處理然后為特定硬件生成可執(zhí)行文件。跟Inductor偏向于“為單個(gè)GPU卡上的小規(guī)模融合”相比XLA的優(yōu)化更傾向于整圖級(jí)別的改寫(xiě)。我最早接觸XLA是在原生TensorFlow時(shí)代那時(shí)候用XLA跑Transformer類(lèi)模型速度提升經(jīng)常是成倍的。后來(lái)PyTorch生態(tài)繁榮起來(lái)torch_xla項(xiàng)目把這種能力搬到了PyTorch里。如果你有TPU資源PyTorch模型可以直接通過(guò)XLA后端在TPU上運(yùn)行這一點(diǎn)是torch.compile目前完全覆蓋不到的。但XLA也有一些明顯的代價(jià)。最大的問(wèn)題是編譯時(shí)間。之前用XLA跑一個(gè)較大規(guī)模的BERT模型編譯階段可能就要幾分鐘到十幾分鐘。如果是訓(xùn)練任務(wù)每個(gè)Step都生成同樣的計(jì)算圖編譯一次就夠了這個(gè)成本可以接受如果是推理任務(wù)每次請(qǐng)求都要處理動(dòng)態(tài)Shape或者新的模型結(jié)構(gòu)編譯開(kāi)銷(xiāo)就會(huì)成為很大的負(fù)擔(dān)。這也是為什么XLA更適用于“固定圖、長(zhǎng)時(shí)間反復(fù)執(zhí)行”的場(chǎng)景。4.2 torch.compile 和 XLA 怎么選很多剛接觸這兩個(gè)概念的人會(huì)糾結(jié)“到底用哪一個(gè)”。我的判斷標(biāo)準(zhǔn)非常簡(jiǎn)單你的目標(biāo)硬件是什么如果你的執(zhí)行環(huán)境是NVIDIA GPU并且模型在PyTorch生態(tài)內(nèi)能跑通默認(rèn)選擇torch.compile。它和PyTorch的接口耦合更緊密調(diào)試信息和工具鏈也更成熟。如果你的目標(biāo)是TPU或者你的模型是從TensorFlow/JAX遷移過(guò)來(lái)的那XLA就是理所當(dāng)然的選擇。另外有些場(chǎng)景會(huì)同時(shí)用到兩者。比如在NVIDIA GPU上做模型開(kāi)發(fā)驗(yàn)證然后跑到TPU上做大規(guī)模訓(xùn)練。我的經(jīng)驗(yàn)是先各自跑通最小實(shí)驗(yàn)再統(tǒng)一對(duì)比指標(biāo)。快速看一張我整理的對(duì)比表對(duì)比維度torch.compile InductorXLA 后端目標(biāo)硬件主要面向NVIDIA GPUNVIDIA GPU / TPU / CPU接入成本一行代碼需要切換設(shè)備為XLA并處理數(shù)據(jù)搬運(yùn)編譯時(shí)間通常幾十秒到幾分鐘大模型可能較長(zhǎng)動(dòng)態(tài)Shape支持較好但允許一定回退一般盡量避免調(diào)試體驗(yàn)較好支持回退Eager模式相對(duì)繁瑣報(bào)錯(cuò)和符號(hào)化程度較高生態(tài)現(xiàn)狀PyTorch官方主推在TensorFlow/JAX場(chǎng)景下更普遍這張表不是絕對(duì)的但它能幫你快速判斷自己該往哪個(gè)方向投入。實(shí)際上我見(jiàn)過(guò)不少團(tuán)隊(duì)為了“追趕熱點(diǎn)”硬上XLA結(jié)果模型在GPU上反而更慢了。因?yàn)閄LA在GPU上的優(yōu)化效果很多時(shí)候并不比Inductor更優(yōu)反而因?yàn)檎麍D編譯的調(diào)度開(kāi)銷(xiāo)在小模型和短任務(wù)上半點(diǎn)便宜都占不到。4.3 XLA 使用中的常見(jiàn)坑位使用torch_xla的時(shí)候最容易踩的坑是數(shù)據(jù)類(lèi)型和Shape不一致帶來(lái)的額外編譯。XLA不喜歡動(dòng)態(tài)Shape。同一個(gè)模型如果每次傳入的sequence長(zhǎng)度都不同XLA就會(huì)反復(fù)做編譯或選擇“動(dòng)態(tài)Shape路徑”性能大打折扣。解決辦法是在數(shù)據(jù)加載階段做好padding把序列長(zhǎng)度統(tǒng)一到同一個(gè)batch內(nèi)的最大值。另一個(gè)坑是分布式訓(xùn)練時(shí)的數(shù)據(jù)同步。torch_xla有自己的分布式接口直接從原生PyTorch的DistributedDataParallel遷移過(guò)來(lái)可能會(huì)碰到collective通信實(shí)現(xiàn)不一致的問(wèn)題。如果只是小規(guī)模實(shí)驗(yàn)單卡訓(xùn)練問(wèn)題不大一旦上多卡建議按照torch_xla官方文檔里的分布式示例調(diào)整代碼而不是硬套原來(lái)的DDP邏輯。最后提一句torch_xla的版本更新頻率通常滯后于PyTorch主版本。如果你正在用很新的PyTorch nightly版裝torch_xla時(shí)最好看一下官方兼容矩陣避免出現(xiàn)API對(duì)不上的問(wèn)題。我自己就遇到過(guò)“小版本不兼容模型A能跑B不能跑”的怪問(wèn)題最終排查出來(lái)是編譯版本不一致導(dǎo)致的。5. 實(shí)操記錄從Eager到torch.compile的一次完整調(diào)優(yōu)這里我會(huì)用一個(gè)簡(jiǎn)化但完整的例子展示我實(shí)際調(diào)優(yōu)一個(gè)CV識(shí)別模型的步驟。這個(gè)模型不復(fù)雜但足以說(shuō)明編譯優(yōu)化的整體流程和注意點(diǎn)。你完全可以把這套流程套用到自己的模型上。5.1 第一步先跑通Baseline拿到可量化指標(biāo)任何優(yōu)化工作第一步永遠(yuǎn)是拿到準(zhǔn)確、可復(fù)現(xiàn)的Benchmark基線。不要上來(lái)就改代碼加編譯否則優(yōu)化前后對(duì)比會(huì)出現(xiàn)很大的噪音。我通常固定隨機(jī)種子、固定輸入Tensor的Shape、固定優(yōu)化器參數(shù)并且把Warmup步數(shù)留足再統(tǒng)計(jì)穩(wěn)定的Step時(shí)間。下面是簡(jiǎn)化示例import torch import torch.nn as nn import time class SimpleModel(nn.Module): def __init__(self): super().__init__() self.conv1 nn.Conv2d(3, 64, kernel_size3, padding1) self.bn1 nn.BatchNorm2d(64) self.conv2 nn.Conv2d(64, 128, kernel_size3, padding1) self.bn2 nn.BatchNorm2d(128) self.pool nn.AdaptiveAvgPool2d((1, 1)) self.fc nn.Linear(128, 10) def forward(self, x): x torch.relu(self.bn1(self.conv1(x))) x torch.relu(self.bn2(self.conv2(x))) x self.pool(x).flatten(1) x self.fc(x) return x model SimpleModel().cuda().train() optimizer torch.optim.SGD(model.parameters(), lr0.01) x torch.randn(32, 3, 224, 224).cuda() target torch.randint(0, 10, (32,)).cuda() criterion nn.CrossEntropyLoss() # warmup for _ in range(10): optimizer.zero_grad() loss criterion(model(x), target) loss.backward() optimizer.step() # benchmark: 50 iters torch.cuda.synchronize() start time.time() for _ in range(50): optimizer.zero_grad() loss criterion(model(x), target) loss.backward() optimizer.step() torch.cuda.synchronize() avg_step (time.time() - start) / 50 print(fEager avg step: {avg_step * 1000:.2f} ms)跑這類(lèi)腳本時(shí)注意每次循環(huán)都執(zhí)行torch.cuda.synchronize()否則計(jì)時(shí)結(jié)果會(huì)不準(zhǔn)確。GPU執(zhí)行是異步的CPU端的time.time()只能記錄到“提交任務(wù)”的時(shí)間而不是真實(shí)執(zhí)行結(jié)束的時(shí)間。5.2 第二步開(kāi)啟torch.compile并觀察Graph Break在Baseline跑通以后第二步就是一行接入compile并觀察Dynamo的捕獲情況。代碼改動(dòng)很簡(jiǎn)單model torch.compile(SimpleModel().cuda().train(), modereduce-overhead)跑同樣的Benchmark流程。如果代碼沒(méi)有特殊控制流torch.compile一般能直接接管大部分計(jì)算路徑。為了確認(rèn)優(yōu)化覆蓋到了多少比例我強(qiáng)烈建議先跑一次帶有解釋信息的診斷from torch._dynamo import explain compiled_model torch.compile(model, modereduce-overhead) explanation explain(compiled_model, x) print(explanation)你會(huì)看到Graph Break的具體位置、覆蓋算子的比例、以及被優(yōu)化掉的節(jié)點(diǎn)數(shù)。如果Graph Break數(shù)量特別多先不要慌逐條分析原因。常見(jiàn)的原因就那幾類(lèi)調(diào)用了.item()、print到Tensor、使用了不支持的第三方庫(kù)函數(shù)、或者對(duì)Tensor做了Python層面的條件判斷。我當(dāng)時(shí)調(diào)這個(gè)簡(jiǎn)化模型時(shí)最典型的問(wèn)題是在forward里加了幾個(gè)用于調(diào)試的assertDynamo遇到這些Python斷言就Graph Break了。把這些調(diào)試邏輯拿到模型外部之后Graph Break數(shù)量從好幾個(gè)降到零整體性能立刻上了一個(gè)臺(tái)階。5.3 第三步對(duì)照組交叉驗(yàn)證區(qū)分“真實(shí)提升”和“測(cè)量噪音”三步跑完通常得到一張類(lèi)似這樣的表模式平均Step時(shí)間相對(duì)Eager提升Eager18.6 ms基線torch.compile(default)13.8 ms~26%torch.compile(reduce-overhead)12.9 ms~31%torch.compile(max-autotune)12.7 ms~32%這個(gè)例子里reduce-overhead和max-autotune的差距已經(jīng)很小max-autotune的編譯時(shí)間卻可能是后者的幾倍。所以在實(shí)際業(yè)務(wù)里我會(huì)優(yōu)先選擇reduce-overhead因?yàn)樗骖櫫司幾g速度和性能收益。需要提醒的是這類(lèi)數(shù)字受很多因素影響GPU型號(hào)、PyTorch版本、CUDA版本、輸入Shape、Batch Size、是否開(kāi)啟AMP等。同一份代碼在不同機(jī)器上可能跑出完全不同的對(duì)比結(jié)果。我自己有過(guò)一次經(jīng)歷在V100上torch.compile只帶來(lái)5%的提升換到A100上同樣的模型直接提升40%。原因很簡(jiǎn)單——新架構(gòu)對(duì)融合Kernel的執(zhí)行效率更高舊架構(gòu)受限于寄存器、調(diào)度能力收益自然不明顯。5.4 實(shí)戰(zhàn)補(bǔ)遺混合精度和編譯的配合如果你在訓(xùn)練中還開(kāi)啟了AMP自動(dòng)混合精度建議把優(yōu)化順序排成“先AMP再compile”。因?yàn)锳MP本身就能把部分計(jì)算切到FP16/FP16帶來(lái)顯著的顯存和速度收益。此時(shí)再疊加torch.compile還能在融合Kernel里自動(dòng)處理FP16/FP32之間的轉(zhuǎn)換細(xì)節(jié)進(jìn)一步減少讀寫(xiě)開(kāi)銷(xiāo)。有一個(gè)細(xì)節(jié)值得單獨(dú)說(shuō)在Eager模式下AMP會(huì)和torch.cuda.amp.GradScaler配合使用避免梯度下溢。在torch.compile下這個(gè)邏輯本質(zhì)不變但Dynamo會(huì)拿到包含AMP包裝層的計(jì)算圖。正常情況下沒(méi)有問(wèn)題但如果你的模型或損失函數(shù)里有自定義的autograd.Function可能需要額外檢查一下是否兼容編譯路徑。出現(xiàn)問(wèn)題的時(shí)候日志里會(huì)明確提示某個(gè)算子不支持編譯這時(shí)候可以把那個(gè)算子所在的子模塊排除出編譯范圍比如用torch.compile(model, disable...)或者把模塊內(nèi)部改成torch._dynamo.mark_dynamic之類(lèi)的注解去處理。6. 踩坑筆記Graph Break、CUDA Graph 與編譯時(shí)間失控理論聊得再多實(shí)踐中該踩的坑一個(gè)都不會(huì)少。這一節(jié)我把自己在torch.compile、Triton和XLA身上踩過(guò)的高頻坑集中列出來(lái)。每一類(lèi)坑背后都是我曾花掉的幾個(gè)小時(shí)能讓你避開(kāi)的話這篇文章就值回票價(jià)了。6.1 Graph Break導(dǎo)致的“越優(yōu)化越慢”最迷惑人的問(wèn)題就是“編譯之后反而比Eager更慢”。這種情況九成以上都能追溯到Graph Break過(guò)多或編譯覆蓋范圍過(guò)小。一旦模型在關(guān)鍵循環(huán)里頻繁回退到Eager那么每次切換到編譯區(qū)又要額外付出圖捕獲和編譯檢查的代價(jià)結(jié)果就是“反復(fù)橫跳”性能自然惡化。排查的方法特別簡(jiǎn)單用torch._dynamo.explain把優(yōu)化報(bào)告打印出來(lái)里面會(huì)清晰地列出每個(gè)Graph Break的文件名、行號(hào)和原因。我遇到過(guò)一個(gè)項(xiàng)目模型里用了一個(gè)外部庫(kù)的某個(gè)函數(shù)對(duì)Tensor做掩碼處理這個(gè)函數(shù)內(nèi)部包含Python層面的循環(huán)和條件分支Dynamo拿它沒(méi)辦法整段都變成了Graph Break。后來(lái)我用原生PyTorch算子重寫(xiě)了那段邏輯Graph Break立刻消失整體訓(xùn)練時(shí)間縮短了接近一半。實(shí)際操作中還有一個(gè)比較隱蔽的trigger是torch.no_grad()和model.eval()的組合使用。推理階段在關(guān)閉梯度后Dynamo的捕獲反而可能因?yàn)榘b關(guān)系變得更復(fù)雜。如果遇到推理階段Graph Break異??梢栽囋嚢裯o_grad移到調(diào)用compile的模型之前或者將推理函數(shù)整體包進(jìn)一個(gè)torch.no_grad()裝飾的函數(shù)里。6.2 CUDA Graph和動(dòng)態(tài)Shape的沖突reduce-overhead模式會(huì)啟用CUDA Graph技術(shù)。CUDA Graph的精髓是“先錄制后回放”第一步執(zhí)行時(shí)把所有Kernel的調(diào)用信息記錄下來(lái)之后就用極低的CPU開(kāi)銷(xiāo)反復(fù)提交。但錄制意味著Kernel的Shape和輸入Tensor的指針信息基本固定。如果之后某一步傳入的輸入Shape變了CUDA Graph回放了錯(cuò)誤配置就會(huì)導(dǎo)致崩潰、顯存錯(cuò)誤或嚴(yán)重性能下降。我的經(jīng)驗(yàn)是訓(xùn)練中若Batch Size固定、圖像分辨率固定GPU Graph幾乎無(wú)副作用。但你一旦在訓(xùn)練過(guò)程中改變了輸入分辨率比如圖像縮放、動(dòng)態(tài)裁剪或者模型中依賴(lài)數(shù)據(jù)的Shape計(jì)算比如序列長(zhǎng)度變化就必須謹(jǐn)慎。更穩(wěn)妥的方案是把變長(zhǎng)輸入先整理成固定Shape的Batch盡量把動(dòng)態(tài)變化控制在模型外部。如果實(shí)在無(wú)法避免動(dòng)態(tài)Shape那reduce-overhead模式建議先不要用改用default模式跑通驗(yàn)證再考慮進(jìn)一步調(diào)優(yōu)。還有一個(gè)細(xì)節(jié)在使用CUDA Graph時(shí)如果模型內(nèi)部有隨機(jī)性操作比如Dropout錄制后的回放階段會(huì)把隨機(jī)數(shù)生成路徑也一并固定導(dǎo)致結(jié)果不可復(fù)現(xiàn)或者統(tǒng)計(jì)分布偏移。PyTorch的編譯器會(huì)對(duì)這類(lèi)隨機(jī)操作做特殊標(biāo)注但手動(dòng)寫(xiě)的一些自定義Triton Kernel如果內(nèi)部調(diào)用了隨機(jī)API需要格外小心。最好在自定義Kernel里顯式傳入隨機(jī)狀態(tài)或者在模型外部控制隨機(jī)性。6.3 編譯時(shí)間太長(zhǎng)是硬件問(wèn)題還是模型問(wèn)題編譯時(shí)間很多時(shí)候讓人抓狂。max-autotune模式在大模型上動(dòng)輒十幾分鐘、幾十分鐘這在開(kāi)發(fā)迭代階段會(huì)非常影響效率。我的做法是“開(kāi)發(fā)用default上線用max-autotune”開(kāi)發(fā)階段頻繁改代碼沒(méi)必要花大量時(shí)間等待編譯模型定型、進(jìn)入長(zhǎng)期固定訓(xùn)練或反復(fù)推理時(shí)再逐步升級(jí)到更激進(jìn)的編譯選項(xiàng)。另外還有一個(gè)容易忽視的原因Inductor為每個(gè)編譯的圖生成多個(gè)候選Triton Kernel并逐一套用不同參數(shù)運(yùn)行驗(yàn)證這個(gè)“autotune”過(guò)程需要在GPU上真實(shí)跑一遍。如果你同時(shí)開(kāi)了多個(gè)進(jìn)程做編譯GPU資源會(huì)被爭(zhēng)搶autotune的時(shí)間也會(huì)被明顯放大。建議在關(guān)鍵編譯任務(wù)執(zhí)行時(shí)保持同一張卡上只有一個(gè)編譯進(jìn)程。如果你發(fā)現(xiàn)編譯總是卡在某個(gè)特定階段還可以開(kāi)啟TORCH_LOGSinductor環(huán)境變量查看詳細(xì)的編譯日志定位到底是圖優(yōu)化階段耗時(shí)還是Triton Kernel自動(dòng)調(diào)優(yōu)階段耗時(shí)。拿到日志后再?zèng)Q定是換mode還是調(diào)整Shape設(shè)定會(huì)高效得多。6.4 自定義算子無(wú)法編譯三個(gè)層面的應(yīng)對(duì)順序PyTorch生態(tài)里難免有第三方或自制算子。遇到Dynamo不認(rèn)識(shí)的函數(shù)時(shí)第一選擇是改成原生算子組合第二選擇是給這個(gè)算子注冊(cè)一個(gè)“偽量化”的圖捕獲規(guī)則讓編譯期可以忽略其內(nèi)部細(xì)節(jié)只把它當(dāng)作一個(gè)不可拆分的節(jié)點(diǎn)第三選擇才是回到Eager模式把整個(gè)模型或某個(gè)子模塊排除在編譯范圍外。在實(shí)際項(xiàng)目中我見(jiàn)到最多的是很多人一遇到自定義算子不支持就慌張地把torch.compile全局關(guān)閉。如果因?yàn)閹讉€(gè)點(diǎn)損失整張圖的優(yōu)化潛力非??上?。更好的做法是把自定義算子封裝在獨(dú)立的子模塊里用torch.compile只編譯模型中其余穩(wěn)定的部分。PyTorch本身是支持部分模塊編譯的善用這個(gè)能力很多兼容性問(wèn)題都可以繞過(guò)去。7. 更進(jìn)一步在真實(shí)業(yè)務(wù)里如何把編譯優(yōu)化推向生產(chǎn)環(huán)境如果你試過(guò)了torch.compile指標(biāo)也漂亮接下來(lái)要思考的是“怎么讓它在生產(chǎn)環(huán)境穩(wěn)定跑起來(lái)”。這涉及的環(huán)境問(wèn)題比模型代碼本身更多比如服務(wù)化部署時(shí)的請(qǐng)求Shape變化、多卡并行時(shí)的組網(wǎng)方式、以及Cluster里多個(gè)任務(wù)對(duì)GPU的爭(zhēng)搶。7.1 訓(xùn)練場(chǎng)景怎么穩(wěn)定落地訓(xùn)練場(chǎng)景相對(duì)簡(jiǎn)單一些。因?yàn)橛?xùn)練數(shù)據(jù)的Shape通常固定模型結(jié)構(gòu)也穩(wěn)定torch.compile的收益可以平滑落地。需要額外關(guān)注的是多機(jī)多卡的通信開(kāi)銷(xiāo)。當(dāng)你把編譯后的模型接入DDP或FSDP時(shí)通信模式和Kernel執(zhí)行順序可能會(huì)和普通Eager模型略有差異如果出現(xiàn)卡頓或利用率波動(dòng)可以先關(guān)掉編譯模式做一下對(duì)照實(shí)驗(yàn)判斷問(wèn)題是出在編譯優(yōu)化還是通信調(diào)度。FSDP的數(shù)據(jù)流比較復(fù)雜Dynamo在捕獲時(shí)會(huì)看到很多與通信相關(guān)的集合操作。新版本PyTorch已經(jīng)對(duì)這類(lèi)帶通信操作的計(jì)算圖做了更好的適配但如果你用的是較老版本遇到FSDP torch.compile性能不升反降的情況可以嘗試把模型內(nèi)部的通信包裝層移動(dòng)到編譯區(qū)之外或者用官方文檔推薦的版本組合。7.2 推理場(chǎng)景怎么處理動(dòng)態(tài)負(fù)載推理任務(wù)要面對(duì)的最大變量是請(qǐng)求Shape不固定。線上服務(wù)經(jīng)常一個(gè)請(qǐng)求是短文本下一個(gè)就是長(zhǎng)文本還有各種不同Batch Size的準(zhǔn)備策略。如果每個(gè)Shape都觸發(fā)一次torch.compile編譯開(kāi)銷(xiāo)會(huì)完全吃掉性能收益。所以生產(chǎn)落地時(shí)一般會(huì)做“按Shape緩存編譯結(jié)果”的設(shè)計(jì)只對(duì)常見(jiàn)的幾個(gè)Shape組合做預(yù)編譯其他Shape走Eager路徑。我在一個(gè)B端推理服務(wù)里用的策略是在服務(wù)啟動(dòng)階段用幾個(gè)典型Shape預(yù)熱編譯結(jié)果然后用一個(gè)哈希表記錄“Shape簽名 - 編譯后模型”。請(qǐng)求進(jìn)入時(shí)先查表命中就使用編譯版本未命中則走Eager。這樣既保證了服務(wù)質(zhì)量又把編譯成本限制在可接受的范圍內(nèi)。論文里有一個(gè)詞叫“Shape Bucketing”說(shuō)的基本就是這個(gè)思路——把連續(xù)的Shape空間離散化成若干桶以兼容性能和靈活性。7.3 對(duì)“性能工程”這件事本身的思考在我寫(xiě)這套學(xué)習(xí)筆記的整個(gè)過(guò)程中一個(gè)反復(fù)出現(xiàn)的主題是性能優(yōu)化沒(méi)有銀彈。torch.compile很好用但它只是把計(jì)算圖表示、算子調(diào)度和硬件指令生成之間的一層又一層的復(fù)雜性封裝了起來(lái)。真正的性能工程能力不是在某個(gè)模型上跑通一個(gè)API而是能快速定位瓶頸、判斷優(yōu)化手段的適用邊界以及在效率和穩(wěn)定性之間做出合適的權(quán)衡。這需要一種“分層診斷”的思維先看數(shù)據(jù)加載和CPU側(cè)是否飽和再看GPU Util是否真實(shí)有效再看計(jì)算圖里是否存在大量小Kernel最后才輪到是否引入編譯器。整個(gè)順序反了你很可能花了一整天時(shí)間調(diào)編譯參數(shù)最后發(fā)現(xiàn)瓶頸在DataLoader根本沒(méi)喂飽GPU。8. 番外幾個(gè)你一定會(huì)用到的環(huán)境與版本問(wèn)題因?yàn)檫@套筆記發(fā)布后經(jīng)常有讀者私信問(wèn)我環(huán)境配置的問(wèn)題這里把跟torch.compile、Triton、XLA相關(guān)的環(huán)境適配問(wèn)題單獨(dú)拿出來(lái)說(shuō)一遍免得大家在前面代碼沒(méi)跑起來(lái)的時(shí)候就被環(huán)境卡住。8.1 PyTorch與CUDA的版本匹配不同版本的PyTorch對(duì)CUDA的支持版本不同。我個(gè)人的習(xí)慣是直接去PyTorch官網(wǎng)的Get Started頁(yè)面選擇合適的操作系統(tǒng)、包管理工具和CUDA版本然后復(fù)制對(duì)應(yīng)的安裝命令。不要自己手動(dòng)從一堆索引里拼輪子那樣很容易引入版本沖突。如果你是離線環(huán)境需要先在一臺(tái)聯(lián)網(wǎng)的機(jī)器上把對(duì)應(yīng)的wheel包下載好然后再搬到目標(biāo)機(jī)器安裝。下載時(shí)注意選擇cu118、cu121這類(lèi)后綴確保和機(jī)器上實(shí)際安裝的驅(qū)動(dòng)兼容。驅(qū)動(dòng)本身要符合CUDA運(yùn)行時(shí)的最低版本要求否則即使torch裝好了torch.cuda.is_available()也可能返回False。8.2 如何確認(rèn)Triton已經(jīng)正確安裝和可用最簡(jiǎn)單的方式是在Python環(huán)境里執(zhí)行以下代碼import triton print(triton.__version__)如果打印出版本號(hào)基本就說(shuō)明裝好了。接下來(lái)再跑一個(gè)小Kernel驗(yàn)證確保運(yùn)行路徑也正確。比如官方文檔里的向量加法示例或者一個(gè)最簡(jiǎn)單的torch.compile測(cè)試。只import成功不一定代表能正常編譯Kernel因?yàn)門(mén)riton還依賴(lài)本機(jī)的編譯器工具鏈。在容器化環(huán)境中經(jīng)常出現(xiàn)“宿主機(jī)能跑但容器里不能跑”的情況絕大多數(shù)是因?yàn)槿萜麋R像里缺少libgomp等動(dòng)態(tài)庫(kù)或者LD_LIBRARY_PATH沒(méi)有正確包含CUDA的庫(kù)路徑。遇到這類(lèi)問(wèn)題時(shí)先檢查基礎(chǔ)鏡像是否包含完整的CUDA runtime再檢查nvcc是否可用。8.3 我常用的一個(gè)快速驗(yàn)證腳本最后分享一個(gè)我?guī)缀趺看芜w移環(huán)境后都會(huì)跑的快速驗(yàn)證腳本內(nèi)容很簡(jiǎn)單但它能在十分鐘內(nèi)暴露80%的常見(jiàn)環(huán)境問(wèn)題import torch print(PyTorch version:, torch.__version__) print(CUDA available:, torch.cuda.is_available()) if torch.cuda.is_available(): print(CUDA version:, torch.version.cuda) print(GPU name:, torch.cuda.get_device_name(0)) x torch.randn(1000, 1000, devicecuda) y torch.matmul(x, x) print(matmul ok:, y.shape) try: import triton print(Triton version:, triton.__version__) except ImportError as e: print(Triton import error:, e) def simple_add(a, b): return a b compiled_add torch.compile(simple_add) out compiled_add(x, x) print(torch.compile ok:, out.shape)如果這個(gè)腳本全綠那說(shuō)明環(huán)境基本可用如果哪一步紅了就跟著堆棧信息去查對(duì)應(yīng)依賴(lài)。每次跑到全綠我才會(huì)開(kāi)始正式的模型優(yōu)化工作。這套流程幫我省去了無(wú)數(shù)次“改了模型卻不知道是環(huán)境問(wèn)題還是代碼問(wèn)題”的無(wú)謂排查。這篇筆記從torch.compile的內(nèi)部機(jī)制講到Triton如何生成GPU Kernel再對(duì)比了XLA這條跨硬件編譯器路線最后落到工程環(huán)境里的落地實(shí)踐。這些內(nèi)容都是我實(shí)際跑模型、調(diào)GPU、被各種版本和Graph Break折騰完之后沉淀出來(lái)的。按我自己的經(jīng)驗(yàn)真正理解這三條主線之后你再去看PyTorch相關(guān)的性能優(yōu)化問(wèn)題視角會(huì)和以前明顯不一樣——不再只是調(diào)幾個(gè)參數(shù)看數(shù)字而是能判斷瓶頸在哪一層、哪條優(yōu)化路徑最適合當(dāng)前場(chǎng)景、換了硬件后又會(huì)發(fā)生什么變化。這一篇先寫(xiě)到這第十五篇打算展開(kāi)聊聊分布式訓(xùn)練里通信和計(jì)算的流水線重疊那個(gè)方向也是我們生產(chǎn)環(huán)境里吃了不少苦頭才摸清門(mén)道的話題。
返回列表
PREV
查看更多資訊
NEXT
返回資訊列表
成人免费看吃奶视频网站| www.91视频网| 日产狠狠干| 午夜天堂精品久久久久91| 中国AAAAAA黄色片| 蜜臀va69| 精品少妇一区二区三区| wwe 天天干.com| 欧美性爱精品一区二区| 国产高清成人mv在线观看| 亚洲天堂另类| 亚洲脚交| 98福利在线视频| 亚洲老熟妇xxx| 日本一天色道久久久精品视频| 91GD.COM| 色色色色网站| 最新啪啪视频| 在线不卡视频| HEYZO高无码国产精品227| 97精品一区二区视频| 九九久久国产精品怡红院| 午夜久久一区二区无码中出| 婷婷丁香六月| 欧美亚洲综合色| 黄色片大香蕉| 亚洲日韩成人性爱视频| 无码少妇精品一区二区60岁老人| 一起草日韩| 强奸乱伦AV一天堂网| 91久久| 成人草草视频| 欧美国产婷婷久久| 亚洲一二三四区| 阿姨一区二区免费视频-高清正片西瓜视频下载app-T450AV | 久久九九视频九九视频| 久久黄色视频一区二区三区| 777琪琪午夜免费A片| 精品久久久久,69国产成人精| 操逼操网| 综合久久中文字幕综合日韩精品| 成人av动漫在线观看| 一,爱啪啪,在线免费视频| 第四色奇米影视777| 国产精品黑人一区二区三区| 性在久久久久久| 日本天天操| 成人性爱美曰韩| av凤凰久久久| 蜜臀99久久精品久久久懂爱| 一中国女人毛片水真多| 操迟操逼在巾线Fre看| 日韩有码 一区二区三区| 91精品国产91久久福利| 国产美女高潮叫床视频| 五月综合色| 亚洲综合五月天| 天天综合AV| 亚洲欧美日韩有码| 日韩有码回春沙龙第一页| 男人的天堂三级| 91操熟女视频| 色吧5亚洲| 日韩成人高清一区二区| 亚洲va有码在线天堂| 91欧美性| 99re在线| 屁屁影院一区二区三区国产| 1204金沙人妻懂旧版免费| 91无码人妻| 熟妇色99| 综合 欧美 亚洲 日本| JULIA一区二区三区在线播放| 夜夜操2028| 嗯阿好爽好紧| 欧美综合自拍亚洲综合图| 婷婷五月影院| 欧美激情亚洲情色| 天天躁日日躁狠狠躁| 九九视频黄色片| 黑人精品成人一区二区三区| 91精品国产综合久久久蜜臀| 亚洲天堂久久| 天天爽夜夜爽夜夜爽精| 久久久久久午夜男人的天堂| 东方亚洲在线操逼天堂| 91老熟女视频| 九九热AV| 伊人9| 精品人妻中文字幕高清| 国产精品熟女乱伦| 亚洲大胆人体av| 蜜臀久久99精品久久久久久婷婷| 亚洲精品aa久久伊人| 精品人妻一区二区三区-国产| 嗯嗯啊好大| 1024精品在线| 久久这里只精品99re66图| 超硑97精品| 91亚洲网站| 少妇xx精品| 婷婷丁香一区二区三区| 影音先锋乱伦资源| 夜夜高潮夜夜爽高清视频一| 午夜毛片亚洲精品片国产久久久| 日韩性爱再线视频| 青青伊人这里只有精品| 青青草男人天堂| 国产h小视频在线观看免费| 自拍大香蕉乱插| 秋霞午夜视频一区二区| 免费A片三p视频| 男人天堂2019亚洲| 日韩精品人妻| 色妺妺在线视频| 九九九九一区| 99热国产| 国内精品久久久久影院亚洲| 中文字幕第23区| 久久精品国产亚洲AV无码做| 91久久久老司机| 无码精品久久| 99re在线观看| 国产av强奸美女| 啊灬快c我灬啊灬用力灬啊灬-国产精品性做久久久久久-成人AV | 九九久久国产精品怡红院| 伊人久久综合影院| 99re这里| 久久久 国产精品| 亚洲综合骚逼| 久久久夜夜嗨免费视频| 9I1性色影院| 色色97爱| 久久精品免视看国产成人﹣蜜臀av一区. 久久精品免视看国产成人,蜜臀av一区 | 熟妇熟女一区二区三区| 亚洲18禁| 99亚洲精品| 97色色色综合网站| 不卡免费av在线播放| 另类老少妇| 凹凸视频特色日本特黄| 偷拍视频青青草在线视频| 99色色| 人妻少妇精品久久久久久久| n1038 一二三区| 日本一区视频在线观看| 精品一区二区2| 色色激情五月天| 超碰人妻久久| 青青草大香蕉在线视频| 日韩少妇一区二区三区| 久久久久免费少妇| 久久中文字幕一区不卡| 嗯~啊~快点 死我视频免费看网站| 黄色网址在线免费观看| 午夜舔阴达高潮视频免费看| 久久国产999| 在线另类| 亚洲欧美综合区自拍另类| 欧美一区二区三熟女剧情| 色老汉玖玖爱| 亚洲日韩美国人妻| 色噜噜国产精品视频一区二区| 日韩专区数据列表-第3230页-精品国产一区二区三区香蕉 久久99熟女人妻中文字 | 丁香激情五月| 日本十八禁免费看污网站| 破处bbq| 婷婷五月天补不补| 在线小说视频一区| 殴美日韩m| 人人干人人操人人..com| 日夜精品| 九九久久国产精品怡红院| 亚州伊人色综台| 亚洲无码?第一页| 国产精品久久久久久久久久久久久久吹 | 熟女这里只有精品6| 超碰在线综合97| 观看免费区二区三区二| 国产激情在线观看| 大香蕉黄色一级片免费看| 9Ⅰ超碰| 久草精品视频| 日日橹狠狠爱欧美超碰| 91原创在线观看| 免费看欧美美女黄色大片| 观看免费区二区三区二| 色婷婷亚洲婷婷| 欧美日韩免费性爱| 男女日B国产| 一区二区三区麻豆| 亚洲欧美骚| 色五月AV在线| 97超视频在线观看| 青青草伊人久久| 久久久婷| 久操热线| 国产欧美一区激情交| 秋霞久久亚洲精品成人| 被窝影院午夜看片无码| 国产强奸乱伦xd| 精品无码一区二区三区色欲| 97极品无码| 97在线播放 | 91N综合网| 欧美大香蕉同搞| 操逼天美3区| 国产一级高清免费观看| 欧美亚洲| 超碰97人妻免费在线| 97伊人网| 超碰日韩人妻| www.久久制服糖| 欧美片第一页| 亚洲中文字幕精品一区| 国产无码久久高清| 天天干天天舔| 翔田千里AⅤHD无码| 豆花视频操逼网址| 蜜乳Av成人片网站| 成人性爱美曰韩| 成人免费福利网站国产| 日本123区操B视频| 99这里有精品视频| 亚州情色j区| 欧美黄页在线| 大香蕉免费3| 涩涩涩综合| 精品国产91av一区二区三区| 最近2019中文字幕国语免费版| 探花在线免费观看视频国产一区| 人人插人人搞人人操| 日韩人妻无码专区| 亚洲高清内射| 日本最新1区2区3区| 亚洲精品黑丝| 亚洲无992tv| 色网在线视频观看免费| 色色五月婷| 天天干天天拍| 97网址97| 97天堂| 欧美自拍偷拍综合图片| 国产 无码 一区二区| 成人av免费观看| 久久久蜜桃一区二区三区| 天美av在线观看| 日韩强奸av| 欧美东京热精品A∨| 欧美视频在线视频免费va| 蜜桃臀av在线观看| 久久久久密臀视频| 久久免费少妇| 欧美日韩性爱无码| 精品丝袜无码一区二区三APP| 人妻丰满熟妇一区二区三| 亚洲一卡2卡3卡4卡乱码网站| 欧美中出1| 96AV久久久| 亚洲精品 欧美精品| 18精品一二区| 天天草夜夜草高潮片| 日韩无码成人电影| 亚洲欧洲色情高清| 熟妇色99| 啊啊啊啊一区| 亚州Av天美传媒| 人人弄人人摸| 亚洲图片91| 国产强奸乱伦xd| 成人欧美一区二区三区黑人一| 亚洲污一污二| 97超碰色色| 性在久久久久久| 蜜臀久久一区二区| 综合网亚洲1| www熟女乱伦com| 韩国黄色片精品久久久| 精品日韩产品在线,日韩在线不卡视频,欧美日韩免费专区/久, | 日韩内| 嫩草影院在线观看精品| 97在线欧洲| 啊啊啊想要| 超碰99热| 蜜臀久久99'精品久久久| 人人操人人爽人人操人人| 九九色逼| 色欲天香天天综合网-成年人三级片网站-欧美乱妇狂野-日韩国产专区-久久久久久 | 岛国免费黄色网址| 欧美天堂第二区| a片 xxxx受爽视频| 狠狠综合| 天天亚洲| 蜜桃久久综合视频| 亚洲精品99999| 色情成人五月天| 91精品操美女| 97久久久| 美欧色综合| 午夜天堂啪啪| 日天天九九天堂666| 成人激情无码在线视频| 色婷婷一区二区三区久久午夜| 天天日天天搞天天干| 精品人妻夜夜草| 偷拍精品一区二区三区| 国产又黄又粗的视频| 国产男女无套97| 国产精品懂色tv影视免费观看 | 东京热毛片调教| 婷婷天堂站| 亚洲猛交| 欧美最婬乱婬爆婬性视频 | 欧美久久久| 成人a大片在线观看| 欧美熟女丝袜| 亚洲图片欧美日韩| 一本大道不卡一二三区| 欧美综合色| 日本在线一二| 精品丝袜无码一区二区三APP| 久久大线蕉一区| 操操吧亚洲乱伦视频| 欧美成人综合| 综合久久婷婷| 国产精品不卡高清在线观看| 一区二区三区成人| 欧美综合网A| 亚洲男人天堂网久久| 尤物视频一区| 欧美日韩亚洲天堂| 久久国产对白激情浪潮 | 麻豆一区二区三区精品| 超碰成人国产| 婷婷超| 日影院久久婷婷夜夜网| 激情小说五月天| 日韩免费一级性爱视频| 久久6热精品99视频| www.色婷婷色综合| 欧美在线永久天堂| 欧美色性情| 熟妇人妻一区二区三在线 | 蜜臀久久99精品久久久| 久久久精品久久| 亚洲色图超碰在线| 丁香九月激情啪| 日韩免费在线观看不卡| 九九九不卡| 日韩在线观看三级电影| 久久大香蕉97| 亚洲成人妻日韩在线| 欧美黄色片AAAAA| 69av一区二区三区| 福利视频合集| 91九色丨风韵犹存| 久久久三区二区一区| 欧美内射少妇| 日韩欧美亚洲国产日韩| 无码抄逼网| 精品久久无码午夜福利| 麻豆天美91| 亚洲精品97p| 欧美韩国你懂得在线 | 骚人妻少妇视频| 欧美内射少妇| 日韩av女优在线免费一区| 亚洲高清无码AAA久久久精品| 91精品人| 国产精品一区二区a| 久久久青青草| 欧美综色欧| 久久视频,这里只有精品| 97久久超碰国产网站| 久久精品国产亚洲粉嫩| 精品人妻无码一区二区三区不卡-精品人妻无码一区二区...|精品少妇一区二区三 | 91偷拍欧美亚洲| 91中文在线| 亚洲精品a人片在线观看视| 九九精品99| 99婷婷一区二区| 男人的天堂va| 蜜臀99久久精品久久久懂爱| 搡老女人老熟女91老熟女综合网| 国产成人精品网站| 本道综合精品| 天天综合-91入口| 蜜臀久久久99久久久久 | 综合网欧美在线| 午夜人妻精品综合在线| 亚洲麻豆精品二区三区| 午夜天堂精品久久久久91| 欧美精品1区2区3区| 黄色av片三级三级三级免费看| 在线国产福利网址导航 | 亚洲91综合| 久久↗↗| 青青草色AV| 久久爱97| 久艹日日日| 国产精品扒开腿做爽爽爽视频| 婷婷色在线| 熟女精品一区二区三区| 久久久999日本大片| 成人在线视频二区| 夜夜操91744565| 久久AV无码1区2区3区| 偷拍超碰| 天天综合站| 传媒在线观看一区二区三区| 日韩人妻丝袜中文字幕| 自拍偷拍草一草| 欧美天天干| 婷婷丁香五月激情啪啪| 女沟厕偷窥piss小便| 国产精品第一区第一页| 欧美日韩国产高清在线一二三区 | 亚洲AV无码黄色强奸| 久久九色| 国产乱码精品久久久久久| 碰超人人在线一区二区三区| 久久宗合97| 精精品人妻一区二区三区| 国产精品久久久| 欧美激情性久久久久久| 久久美女福利是上海美女| 欧美午夜精品久久久久久超碰| 国产一区二区三区不卡手机在线| 丝袜狠狠草尤物 91| 亚洲天堂男| 精品少妇人妻av久久免费| 日本片日本片祼观看网站在线看中文版网页在线看 | 亚洲精品aa久久伊人| 激情五月激情综合网| 2025亚洲男人天堂| 亚洲色电影在线| 国产精品亚洲一级av第二区| 熟妇在线视频一区二区| 国产精品午夜福利视频| 亚洲精品久久久久毛片A片拉屎 | 综合网亚洲1| 99精品国产户外露出| 欧美劲爆视频一区二区| 欧美日韩人人早| 精品人妻高清麻豆av| 国产精品久久久久久久电影渣男| 激情综合五月| 91欧美巨乳| 国产精品探花色| 91亚州日韩高清| 中文字幕乱偷人妻久久艾草网| 久久男人的天堂| 亚欧国产无码精品在线| 超碰97在线中文| 亚洲无码偷拍| 婷色五月天| 亚洲色综网| 国产成人午夜视频网址| 97色97好| 综合伊人激情| 欧美色图偷拍另类| 高清国产精品福利网站| 五月综合色| 黑丝制服中文字幕| 久久香蕉综合一本到3atv| 亚州综合AⅤ| 亚洲天堂在线怕怕视频| 伊人久久亚洲中文字幕| 精品少妇999| 伊人天堂在线| 欧美天天综合网| 日本一级一级一级一级| 久久粉色| 精品国产乱码久久久久A| 97精品一区二区视频在线观看| 国产男人又猛又粗又爽| 果冻传媒A片一二三区| 可以在线观看AV的网站| 天天影视网色欲色香| 久久久久久久国产视频| 91N综合网在线| 人摸人人操人| 色狠狠一区二区三区香蕉| 睡产熟女乱伦| 青青草原香蕉日本Ap| www狠狠| 日韩91网站| 欧美|91色综合| av网站免费看| 欧美日韩亚洲天堂网| 亚洲一区二区久久久久| 欧美色综合图片| 欧亚性爱啪啪| 久久亚洲一区女同性恋中文字幕| 天天影视色香色欲| 亚洲无码超碰免费| 亚州色图片在线色| 超碰在线国产| 国产综合色精品在线观看| 亚洲宗合电影| 亚洲欧美日韩精品久久久一区二区| 97精品视频在线播放| 天天天天操| 免费精品国偷自产在线在线| 国产熟女二区| 精品色色| 亚洲AV无码翔田千里网站| nuu12国产麻豆精品| 五月丁香久久| 大香焦A片| 四季av一区二区凹凸精品小说| 嗯嗯啊中文字幕| 日韩熟女三十乱伦| 亚洲清纯唯美| 中文字幕在线观看第二页| 亚洲自拍天堂| 亚洲欧美首页| 中文字幕伊人| 亚洲精品色| 国产人妻天天干精品| 亚洲人成在线放东京热| 997色在线| 欧美美女自慰一区二区三区| 久久精品店| 白丝av| 后入式999| 国产专区路线| 99激情视频| 国产熟女完整版中字| 97欧美日韩| 91青青在线视频| 看黑丝美女操逼青青网站| 国产精品老师| 国产精品嫩草影院免费| 极品销魂美女一区二区| 精品少妇99| 日韩久草| 国产中文福利| 国产性刺激| 免费视频无码| 中文字幕亚洲在线一区| 天天做日日做| 欧美人人曰人人操人人射射| 91老女人| 97天天摸天天爽| 九九九只有精品| 伊人aaa| 国产精品亚洲四五区在线观看| 任你草| 黄网色一区二区三区四区精品| 无码人妻精品一区二区三区99不卡| 久久精品无码专区| 91热| 成人性交午夜免费片| 97久久精品不卡| 久久久久密臀视频| 天天看片麻豆| av东京热男人的天堂| 四虎视频在线观看| 草久久久| 97超碰巨乳| 日日AV加勒比| 成人一区二区三区四区| 91美女在线视频| 天天综和| 亚洲www91| 亚洲色91C| 狠狠躁日日躁夜夜躁A| 成人久久精品| 天天综合麻豆视频| 激情文学亚洲| 亚洲深夜福利| 91丨人妻丨国产丨丝袜| 可免费观看的av毛片中日美韩| 日本色日夜干| 午夜亚洲WWW湿好大| 国产福利视频精品视频| 五十路三级片| 凹凸视频特色日本特黄| 日韩欧美视频青青| 91最新综合| 色超碰综合| 国产欧洲精品亚洲午夜拍精品| 欧美在线干| 日日躁夜夜躁狠狠躁超爽| 人妻一区二区三区熟女| 国产无马在线| 少妇高潮对白在线观看| 激情五月婷| 国产久久视频| 本道在线| 天天澡天天狠天天天做| 国产精品亚洲一级av第二区| 日韩av乱伦| 91逼逼女人91| 囯产精品强| 熟妇女伦乱视频| 日韩字幕一区| 熟女突然公开看18禁影片| 91色综| 麻豆成人av| 97操B| 激情天天视频| 中文激情网| 欧美se亚洲| 亚洲春色欧美| 中文日本免费高清| 日韩高潮一区| 精品成人女人久久| 成全在线观看免费观看| 五月天综合| 久久国产精品熟女人妻| 丝袜综合网| 丰满人妻一区二区三区四区| 中文字幕啊啊啊在线观看视频| 夜夜嗨绯色| 东京热毛片177b2viP| 国产东北女人在线视频| 在线观看AV不卡| 亚洲不卡一| 日韩操p| 先锋精品av色鲁| 日日妻色网| 明星性猛交ⅹxxx乱大交| 婷婷久热| 91少妇通奸网站| 草草草视频在线免费看| 亚洲激情在线| 精久久久| 色婷婷激一区二区三区| 影音综合网| 中文字幕乱亚洲美女精品一区| 91成人在线免费视频| 丝袜综合色图| 人人人干干人人干| 啊啊啊好爽快点啊啊啊嗯嗯| 亚洲成人AB| 人妻日日干| 欧美色日| 日本免费一级AAA大片器| 人妻插插人妻人| 亚洲综合第一页| 激情深爱五月天| 精品少妇一区二区| 国产精品国产拍高清AV| 日夜尻逼网| 黄色视频特级毛片| 另类专区加勒比| 亚洲成a人片在线观看中文!!!| 偷拍超碰| 中国一级特黄大片护士| 欧美日韩婷婷中文| 日韩中文字幕精品一区在线| 草伊人高潮喷水超碰| 日韩电影在线观看网址| 人人手机欧洲亚洲国产人妻| 久久久精品国产亚洲AV无码| 97久久久久久久久久| 一个人免费HD91视频| 玖玖爱伊人玖玖爱| 日韩AV熟女乱伦| 午夜视频久久久| 亚洲另类电影| 女生久久网| 精品人妻中文字幕高清| 韩国一区二区精品亚洲| 日韩欧美亚洲自拍偷拍| 嗯……啊…嗯嗯…啊…好舒服| 蜜臀一区二区三区在线 | 久久噜噜噜精品国产亚洲综合| 欧美黑人精品一区二区| 精品人妻一区二区蜜桃视频| www.色婷婷| 91w欧美| 欧洲一区二区三区免费| 国产精品久久久久亚洲av| 国产精品乱码久久久、久久| 婷婷色导航| 久久久久九九九| 天天综合网1| 秋霞曰韩R级| 超碰人妻久久| 丁香五月天久久精品视频一区二区三区| 日本狂喷奶水在线播放212| 九七人妻在线| 亚洲天堂区| 久久免费99精品久久久久久| 中出在线视频| 亚洲中文制服诱惑| 欧美成人AⅤ大片在线观看| 欧美同性恋 的搜索结果 - 91n| 蜜桃色色网站视频三区| 性爱乱伦视频免费| 超碰超碰超碰超碰的大鸡吧操黑丝袜 | 97香蕉人人乳| www.国产高潮精品| 国产在线综合福利网站| 锕锕好爽 死我在线观看| 97久久精品亚洲| 精品久久久久久中文| 欧综合网| 天天操天天看| 久久这里都是精品| 久久精品国产免费观看99| 加勒比综合网| 蜜乳AV.COM| 欧美视频一区二区三区| 亚洲 欧美 91| 日韩丝袜二区| 一级做a爰片性色毛片久久| 综合色久欲| 蜜桃视频一区二区三区在线观看| 超碰99热| 免费伦费视频在线观看| 一起草三级AV电影在线观看 | 日本精品五区| 亚洲少妇自拍中文字幕懂色| 四虎国产成人精品免费一女五男| 天天干,夜夜爽| 乱欲一区二区| 欧洲亚洲国产综合在线| 欧美九九九九九| 国产精品乱码久久久、久久| 欧美熟女妇同| 国产精品一区二区在钱播放| 99精品热| 黄片在线免费在线观看| 草草影院最新网址| 午夜亚洲WWW湿好大| 欧美性色网| 亚洲欧美激情小说| 麻豆黄站| 97免费视频在线观看视频| 99re9这里只有精品| 97色色色| 色欧洲| 国产熟女一区二区| 久久久天堂| 精品高清牛人盗摄一区二区三区中文字幕A片免费在线观看 | 女性喷水高潮在线观看| 校园春色 欧美| 麻豆av一区二区| 亚洲色婷婷综合久久一区二区三区| 好吊爽好吊爽在线视频,中文字幕精品一区二区日本,国产良妇出轨视频在线观看, | 免费在线黄片视频| 亚洲欧美自拍偷拍| 再深点灬舒服灬太大了添视频| 亚洲成aⅴ人片不卡无码| 日韩日本欧美在线观看| 精品国产av一区二区三区四区入口| 无码人妻一区二区三区免费九色| 91天天综合日韩欧美| 欧美乱欲| av毛片aaaaa免费看| 强被迫伦姧在线观看无码网站| 国产精品不卡av免费在线观看| 国产精品懂色tv影视免费观看| 日本视频在线中文字幕| 91色人妻| 亚州免费啪啪视频| 很很热性爱视频| 97婷婷色| 96AV精品| 躁躁日曰躁2020| B049AV在线播放| 蜜桃臀久久| 99re视频在线播放青草| 天堂无码| 国产久久成人| 亚洲综合中文字幕有码| 中文字幕在线免费观看视频| 青青草视频这里只有精品| 婷婷三区| 久久人妻无码毛片A片麻豆| 精品高清牛人盗摄一区二区三区中文字幕A片免费在线观看 | 中文字幕天天天天天| 欧美在线55555| 色一区二区三区综合| 大香蕉人妻久久| 国产精品人妻无码久久久互動交流| 青娱乐二区免费| 久干9操| 成人五月香网在线| 欧美性爱一区二区三区四区| 亚洲蜜臀精品视频久久| 少妇被c 黄 免费观看| 日韩在线一区高清在线| 久热久| 99re9这里只有精品| 精品9999| 亚洲日韩人妻中文字幕一区| 日本一本道A级黄色毛片试看60分钟| 婷婷大香蕉| 蜜臀99999| 亚洲人妻日日日| 免费观看性欧美一级| 日韩中文字幕二区| 中文字幕av久久爽Av| 大干人妻| 青青操网| 婷婷月色| 日日干夜夜欢| 亚洲国产美女久久久久| 大地资源在线观看中文第二页| 亚洲色图 欧美热图 清纯唯美 另类自拍 | 国产日韩中文字幕欧美| 91免费看一区二区三区| 性色av婷婷久久一区二区点复制| AV高清一区| 嗯嗯不要视频| 久久‘黄片视频| 国产性感在线观看| 97亚洲中文| 操逼网免费无码视频| 青青草手机在线免费观看| 国产精品一区二区三| 色成人Www精品永久观看| 91亚洲青青草原精品1区| 91丝袜人妻| 日韩一级特黄av毛片| 色好看av| 国产精品九九九| 婷婷精品视频| 久久久内射良家| 国产曰批免费观看久久久| 伊人色综合超碰| 欧美日韩1234| 99国产人成精品| 极品出轨视频网站| 欧亚性爱视频免费看| 加勒比伊人综合| 成人无遮挡毛片免费看| 欧美色九九九| 色色毛片| 精品国产污一区二区三区| 熟女精品日韩一区二区三区| 熟女六十路| 国产AB视频| 日韩激情中文字幕有码| 蜜臀无码一区二区| 国产亚洲中文不卡二区| 综合视频91| 日韩乱伦AⅤ| 中文字幕人妻色偷偷久久皮 | 在线人人人人人人精品超| 高清肉丝中文无码| 国产精品爽爽v| 韩国国产欧美情侣视频在线| 久久九九视频九九视频| 丝袜色综合| www.人人cao| 超碰欧美COM| 97人人草| 91少妇| 天堂国产AV| 国产女人高潮嗷嗷嗷叫小说| 亚洲黄色网址| 留下AⅤ黄色片| 夜夜黄| 超91综合网| 91视频精品| 好看的久久不射无码影视影院| 日本丝袜人妻内射| 乱色老一区二区三区的观看方式| 大屁股xxxxx| 亚州男人天堂| 亚洲综合影视| 国语国产操逼伊人AV网| 香蕉国产97| 色香色香欲天天天影视综合网| 日韩欧美性吧婷婷乱伦大香蕉| 伊人网在线观看| 另类 日韩 熟女| 青青青在线高清视频在线一二三四区 | 亚洲精品国产熟女久久久| 91丨九色丨43老版熟女| 操东北女人| 亚州高清色综合| 亚洲中文一区二区三区视频| 久久东京热久久| 国产肏屁眼视频| 日韩中文字幕视频在线观看| 性色亚洲| h色99999| 91色综| av在线一区二区三区| 精品人妻丰满熟妇一区二区三| 久久视网78| 一区在线观看中文字幕| 日本在线观看网址| 呻吟 欧美 日本 中出| 国内91熟女人妻丝袜天天精品视频在线 | 伊人网在线点播| 人妻久热在线| 国产女人91精品嗷嗷嗷嗷| 91综合网| 99色在线| 禁十八久久| 婷婷六月色| 久久毛卡| 人人么人人操| 超碰97人妻自拍| 东京热不卡视频| 亚洲 中文 女同| 亚洲偷拍自拍在线视频| 伊人一区二区在线播放| 在线欧美亚洲| 久久久久久久9最新免费视频观看| 炮色五月| 五月婷丁香| 国产精品自在线发布| 午夜啪| 后入国产| 欧美最婬乱婬爆婬牲视频| 97 国产一区| 亚洲 图片 综合91| 日本一级性爱| 亚洲va有码在线天堂| 一级片视频啪啪| 天天干18禁| 久久久久久久久9| 亚洲综合伊人| 亚洲天堂美臀在线| 啊啊啊好舒服好爽啊啊啊视频| 乱人伦 国语对白:视频直接看| 日韩人妻一区二区精品| 啪啪免费| 亚洲欧美在线观看无码| 在线播放成人高清免费视频| 中文字幕三四五区| 欧亚 另类 久| 国产又大又粗又长视频| 欧美最大综合网| 欧美日韩天堂| 伊人国产视频| 日韩AV中文字幕电影| 日韩亚洲美女一区久久| 另类小说五月天| 五月天色色网站| 久久草在线综合视频| 欧美专区日本专区| 操高情无码| 国产偷拍自拍在线视频| 嫩草一区二区在线观看| 国产精品美女在线一区| 思思热国产在线视频| 日韩中文字幕人妻视频| 95自拍视频在线观看| 欧美亚洲AN| 久草色悠悠在线视频| 99国产在线绯色一区| 宗合情欲网| 国产美脚女优尤物在线观看| 欧美老妇综合网| 殴洲老熟女| 人人操人人射人人干| 日韩国产九九精品一区二区三区毛片| 97爱啪| 人人操人人操人妻人| 猛猛干| 久久超碰国产一区二区三区| 一级日本牲交大片好爽在线看| 亚洲精品国产熟女| 偷拍自拍在线视频观看| 日韩免费性爱视频在线观看| AA丁香综合激情| 久久久久久夜夜夜夜夜| 精品国产99| 物业黑人 AV一区| 欧美激情 亚洲色图| 舔足天天操天天射| 亚洲97成人在线观看| 亚洲天堂AV在线播放| 爱媛媛久久国产福利| 夜夜爽33333| 免费人人搞97| 激情文学小说一区二区 | 人人看人人摸人人色| 婷婷综合五月天| 熟妇视频一区二区三区在线观看| 操死我了嗯嗯嗯| 天天操福利视频综合网站| av中亚| 国产专区第一页| 欧美亚洲日韩人妻在线观看| 91在线免费观看处女| 久久国语| 777超碰| 逼操网站| 亚洲中文字幕97久久精品少妇| 国产天天看| 欧美三级偷拍| 九九精品无码专区免费| 国产超碰| 男人的天堂欧美| 国产精品乱码久久久久久久久| 欧美传媒| 色播丁香| 精品人妻无码一区二区三区不卡-精品人妻无码一区二区...|精品少妇一区二区三 | 男人把坤坤插入女人的下体 | 精品四五区| 精品人妻视频一区二区在线播放| 可免费观看的av毛片中日美韩| 亚洲av影院在线观看| 亚熟在线| 999久久久国产精品| 操逼不卡中文字幕| 亚洲不卡不卡中文字幕不卡 | 久久在肏| 国产1769在线| 91撸色网 玖玖网 欧美| 91性网| 日韩欧美午夜视频在线| 色女99一级片在线观看| 老司机香蕉久久久久| 51一区二区三区| 97se综合网| 夜夜爽夜夜| 1024人妻| 在线电影亚洲色图| 欧美天天干| 欧美性暴力猛交XXXX| 亚洲欧美情色| · —级AA伦aa坐爱午夜极速ⅴA一区天天噪天天噪天天噪 | 伊人991| 国产无码三级视频在线观看| 激情文学网伊人| 超碰诱惑| 91丨九色丨东北熟女| 天天爱综合网| 亚洲暴力强奸AV| 天美传媒Av在线| 91av一区二区在线观看| 欧美最婬乱婬爆婬性视频 | 开心五月激情网| 欧美淫乱视频| 人人操人人摸人| 天天热精品| 亚洲美女自拍偷拍视频| 欧美人人天天网| 新久久AV| 精品无码久久久久| av天堂天堂av日韩| 91/欧美| 日日操丁香五月天| 日韩精品电影| 网页导航五月天免费一二三区 | 99热伊人| 天天综合网~91| 97超碰国产亚洲精品资源| 啊啊啊好舒服视频| 粉嫩国产精品久久粉嫩| 久久久97| 亚洲蜜臀懂色| 欧美亚洲一级在线观看| 91强在线播放| 天天干天天狼在线视频| 青青草在线视频播放器| 色999亚洲人成色| 清纯唯美亚洲综合| 久久色一区| 天天懆天天日| 超碰97欧美日韩| 综合色啪| 国产剧情在线| 殴美综合色88| 久操凹凸视频| 成人草草视频| 综合影院亚洲| 天天操人人操狠狠插| 九九九九97| 久久人妻办公室视频| 亚洲欧美日韩免费观看| 97超碰在线资源网站| 亚洲国产欧美中文永久| 男人天堂欧美| 亚洲少妇激情视频| 日本亚洲熟女视频| 黄骗免费网站| 夜草欧美| 久99热| 高清无码久操视频| 亚洲男人天堂视频| 欧美在线官网| A V视频日本| 欧洲亚洲少妇| 日本东京热大香蕉a片| 亚洲高清无码免费观看视频| 人人操,人人液| 久久精品视频一区三区小泽玛利亚| 超碰在线成人| 懂色中文一区二区三区| 人人澡人人澡人人| 国产精品乱人伊人网| 国产成人免费观看在线视频| 91五十路| 综合性视频99| 超碰国产情侣自拍网| 在线观看黄色电话| 伊人网免费视频| 校园春色美腿丝袜 | 99日视频在线免费| 天天亚洲综合| 五月婷丁香| 国产成久久综合片| 超碰日本97美女人妻人人玩人人爱| 午夜操逼不卡| 我要色综合网| 天天操狠狠日夜夜干超大胆开放com大香蕉视频在线观看 | 性爱边摸边日免费AV| 中文字幕在线观看网址| 97在线欧| 日本东京热大香蕉a片| 蜜桃久久综合视频| 人妻99p| 精品毛片av一区二区| 91天天| 91中出视频| 99久久9| 啊啊啊啊二区好大| 校园春色之综合网| 十八禁视频网站| 欧美操逼熟女| 亚洲av国产av综合av卡| 久久九色| sewuyueav| 玖色av| 97欧美久久久久久久| 亚洲中文电影| 亚洲限制级在线| 爽爽淫人网| 啪啪一区| 亚洲天堂 视频你懂的| 亚洲成人无码影院| 精品无码久久久久久国产浪潮| 国产强奸乱伦第1页| 噜噜噜噜久久久精品免费| 欧美日韩97在线| 天天拍天| 青青伊人这里只有精品| 国产少妇高潮| 91五月天| 免费成人在线熟妇网| 久久中日麻豆| 亚洲欧洲无码97久久精品| 亚洲AO在线| 欧美色图天堂网m| 青青色综合| 亚洲色图加勒比| 色综合加勒比四四季| 二对二中文字幕。| 蜜桃不卡一区二区| 天美欧美国产| 蜜桃丰满熟妇av无码区不卡| 欧美亚洲宗合色性图| 少妇第一页| 久久久久幕乱码| 78精品在线| 91 丝袜在线播放| 夜夜骑夜夜操|