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

0
  • 聊天消息
  • 系統消息
  • 評論與回復
登錄后你可以
  • 下載海量資料
  • 學習在線課程
  • 觀看技術視頻
  • 寫文章/發帖/加入社區
會員中心
創作中心

完善資料讓更多小伙伴認識你,還能領取20積分哦,立即完善>

3天內不再提示

PyTorch如何實現自定義CUDA算子并調用的方法且測量CUDA程序耗時

深度學習自然語言處理 ? 來源:算法碼上來 ? 作者:算法碼上來 ? 2021-03-30 15:58 ? 次閱讀

最近因為工作需要,學習了一波CUDA。這里簡單記錄一下PyTorch自定義CUDA算子的方法,寫了一個非常簡單的example,再介紹一下正確的PyTorch中CUDA運行時間分析方法。

完整流程

下面我們就來詳細了解一下PyTorch是如何調用自定義的CUDA算子的。

首先我們可以看到有四個代碼文件:

main.py,這是python入口,也就是你平時寫模型的地方。

add2.cpp,這是torch和CUDA連接的地方,將CUDA程序封裝成了python可以調用的庫。

add2.h,CUDA函數聲明。

add2.cu,CUDA函數實現。

然后逐個文件看一下是怎么調用的。

CUDA算子實現

首先最簡單的當屬add2.h和add2.cu,這就是普通的CUDA實現。

void launch_add2(float *c,

const float *a,

const float *b,

int n);

__global__ void add2_kernel(float* c,

const float* a,

const float* b,

int n) {

for (int i = blockIdx.x * blockDim.x + threadIdx.x;

i 《 n; i += gridDim.x * blockDim.x) {

c[i] = a[i] + b[i];

}

}

void launch_add2(float* c,

const float* a,

const float* b,

int n) {

dim3 grid((n + 1023) / 1024);

dim3 block(1024);

add2_kernel《《《grid, block》》》(c, a, b, n);

}

這里實現的功能是兩個長度為的tensor相加,每個block有1024個線程,一共有個block。具體CUDA細節就不講了,本文重點不在于這個。

add2_kernel是kernel函數,運行在GPU端的。而launch_add2是CPU端的執行函數,調用kernel。注意它是異步的,調用完之后控制權立刻返回給CPU,所以之后計算時間的時候要格外小心,很容易只統計到調用的時間。

Torch C++封裝

這里涉及到的是add2.cpp,這個文件主要功能是提供一個PyTorch可以調用的接口

#include 《torch/extension.h》

#include “add2.h”

void torch_launch_add2(torch::Tensor &c,

const torch::Tensor &a,

const torch::Tensor &b,

int n) {

launch_add2((float *)c.data_ptr(),

(const float *)a.data_ptr(),

(const float *)b.data_ptr(),

n);

}

PYBIND11_MODULE(TORCH_EXTENSION_NAME, m) {

m.def(“torch_launch_add2”,

&torch_launch_add2,

“add2 kernel warpper”);

}

torch_launch_add2函數傳入的是C++版本的torch tensor,然后轉換成C++指針數組,調用CUDA函數launch_add2來執行核函數。

這里用pybind11來對torch_launch_add2函數進行封裝,然后用cmake編譯就可以產生python可以調用的.so庫。但是我們這里不直接手動cmake編譯,具體方法看下面的章節。

Python調用

最后就是python層面,也就是我們用戶編寫代碼去調用上面生成的庫了。

import time

import numpy as np

import torch

from torch.utils.cpp_extension import load

cuda_module = load(name=“add2”,

sources=[“add2.cpp”, “add2.cu”],

verbose=True)

# c = a + b (shape: [n])

n = 1024 * 1024

a = torch.rand(n, device=“cuda:0”)

b = torch.rand(n, device=“cuda:0”)

cuda_c = torch.rand(n, device=“cuda:0”)

ntest = 10

def show_time(func):

times = list()

res = list()

# GPU warm up

for _ in range(10):

func()

for _ in range(ntest):

# sync the threads to get accurate cuda running time

torch.cuda.synchronize(device=“cuda:0”)

start_time = time.time()

r = func()

torch.cuda.synchronize(device=“cuda:0”)

end_time = time.time()

times.append((end_time-start_time)*1e6)

res.append(r)

return times, res

def run_cuda():

cuda_module.torch_launch_add2(cuda_c, a, b, n)

return cuda_c

def run_torch():

# return None to avoid intermediate GPU memory application

# for accurate time statistics

a + b

return None

print(“Running cuda.。。”)

cuda_time, _ = show_time(run_cuda)

print(“Cuda time: {:.3f}us”.format(np.mean(cuda_time)))

print(“Running torch.。。”)

torch_time, _ = show_time(run_torch)

print(“Torch time: {:.3f}us”.format(np.mean(torch_time)))

這里6-8行的torch.utils.cpp_extension.load函數就是用來自動編譯上面的幾個cpp和cu文件的。最主要的就是sources參數,指定了需要編譯的文件列表。然后就可以通過cuda_module.torch_launch_add2,也就是我們封裝好的接口來進行調用。

接下來的代碼就隨心所欲了,這里簡單寫了一個測量運行時間,對比和torch速度的代碼,這部分留著下一章節講解。

總結一下,主要分為三個模塊:

先編寫CUDA算子和對應的調用函數。

然后編寫torch cpp函數建立PyTorch和CUDA之間的聯系,用pybind11封裝。

最后用PyTorch的cpp擴展庫進行編譯和調用。

運行時間分析

我們知道,CUDA kernel函數是異步的,所以不能直接在CUDA函數兩端加上time.time()測試時間,這樣測出來的只是調用CUDA api的時間,不包括GPU端運行的時間。

所以我們要加上線程同步函數,等待kernel中所有線程全部執行完畢再執行CPU端后續指令。這里我們將同步指令加在了python端,用的是torch.cuda.synchronize函數。

具體來說就是形如下面代碼:

torch.cuda.synchronize()

start_time = time.time()

func()

torch.cuda.synchronize()

end_time = time.time()

其中第一次同步是為了防止前面的代碼中有未同步還在GPU端運行的指令,第二次同步就是為了等fun()所有線程執行完畢后再統計時間。

這里我們torch和cuda分別執行10次看看平均時間,此外執行前需要先執行10次做一下warm up,讓GPU達到正常狀態。

我們分別測試四種情況,分別是:

兩次同步

第一次同步,第二次不同步

第一次不同步,第二次同步

兩次不同步

這里我們采用英偉達的Nsight Systems來可視化運行的每個時刻指令執行的情況。

安裝命令為:

sudo apt install nsight-systems

然后在運行python代碼時,在命令前面加上nsys profile就行了:

nsys profile python3 main.py

然后就會生成report1.qdstrm和report1.sqlite兩個文件,將report1.qdstrm轉換為report1.qdrep文件:

QdstrmImporter -i report1.qdstrm

最后將生成的report1.qdrep文件用Nsight Systems軟件打開,我這里是mac系統。

兩次同步

這是正確的統計時間的方法,我們打開Nsight Systems,放大kernel運行那一段可以看到下圖:

0256c144-8e8f-11eb-8b86-12bb97331649.png

其中第1和第3個框分別是cuda和torch的GPU warm up過程,這部分沒有進行線程同步(上面的黃色塊)。

而第2和第4個框就分別是cuda和torch的加法執行過程了,我們可以放大來看看。

02cd92ec-8e8f-11eb-8b86-12bb97331649.png

可以看出,每執行一次(一個框)都經過了三個步驟:先是調用api(左上角藍色框),然后執行kernel(下方藍色框),最后線程同步(右上角黃色框)。

所以最后算出來的時間就是這三個步驟的耗時,也就是下圖選中的范圍:

032b61ce-8e8f-11eb-8b86-12bb97331649.png

時間大概在29us左右,和我們實際代碼測出來的也是比較接近的:

039a9be8-8e8f-11eb-8b86-12bb97331649.png

其實我們實際想要知道的耗時并不包括api調用和線程同步的時間,但是這部分時間在python端不好去掉,所以就加上了。

第一次同步,第二次不同步

放大每次執行的過程:

可以看出,雖然長的和上一種情況幾乎一模一樣,但是在api調用完之后,立刻就進行計時了,所以耗時只有8us左右,實際測出來情況也是這樣的:

047e113e-8e8f-11eb-8b86-12bb97331649.png

第一次不同步,第二次同步

我們先來看一下實際統計的時間:

04eba01e-8e8f-11eb-8b86-12bb97331649.png

很奇怪是不是,第一次運行耗時非常久,那我們可視化看看到底怎么回事:

055a53ec-8e8f-11eb-8b86-12bb97331649.png

可以看出,因為第一次開始計時前沒有同步線程,所以在GPU warm up調用api完畢后,第一次cuda kernel調用就開始了。然后一直等到warm up執行完畢,才開始執行第一次cuda kernel,然后是線程同步,結束后才結束計時。這個過程非常長,差不多有130us左右。然后第二次開始執行就很正常了,因為kernel結束的同步相當于是下一次執行之前的同步。

兩次不同步

先來看看執行情況:

05ef66a8-8e8f-11eb-8b86-12bb97331649.png

可以看出因為沒有任何同步,所有GPU warm up和cuda kernel的api調用全接在一起了,執行也是。所以計時只計算到了每個api調用的時間,差不多在7us左右。

上面四種情況,torch指令情形幾乎一樣,因此不再贅述。

小結

通過這篇文章,應該可以大致了解PyTorch實現自定義CUDA算子并調用的方法,也能知道怎么正確的測量CUDA程序的耗時。

當然還有一些內容留作今后講解,比如如何實現PyTorch神經網絡的自定義前向和反向傳播CUDA算子、如何用TensorFlow調用CUDA算子等等。
編輯:lyn

聲明:本文內容及配圖由入駐作者撰寫或者入駐合作網站授權轉載。文章觀點僅代表作者本人,不代表電子發燒友網立場。文章及其配圖僅供工程師學習之用,如有內容侵權或者其他違規問題,請聯系本站處理。 舉報投訴
  • python
    +關注

    關注

    56

    文章

    4792

    瀏覽量

    84628
  • CUDA
    +關注

    關注

    0

    文章

    121

    瀏覽量

    13620
  • pytorch
    +關注

    關注

    2

    文章

    807

    瀏覽量

    13201

原文標題:【進階】PyTorch自定義CUDA算子教程與運行時間分析

文章出處:【微信號:zenRRan,微信公眾號:深度學習自然語言處理】歡迎添加關注!文章轉載請注明出處。

收藏 人收藏

    評論

    相關推薦

    TPS659xx應用程序自定義工具

    電子發燒友網站提供《TPS659xx應用程序自定義工具.pdf》資料免費下載
    發表于 11-06 10:02 ?0次下載
    TPS659xx應用<b class='flag-5'>程序</b><b class='flag-5'>自定義</b>工具

    創建自定義的基于閃存的引導加載程序(BSL)

    電子發燒友網站提供《創建自定義的基于閃存的引導加載程序(BSL).pdf》資料免費下載
    發表于 09-19 10:50 ?0次下載
    創建<b class='flag-5'>自定義</b>的基于閃存的引導加載<b class='flag-5'>程序</b>(BSL)

    EtherCAT運動控制器PT/PVT實現用戶自定義軌跡規劃

    EtherCAT運動控制器PT/PVT實現用戶自定義軌跡規劃。
    的頭像 發表于 08-15 11:49 ?621次閱讀
    EtherCAT運動控制器PT/PVT<b class='flag-5'>實現</b>用戶<b class='flag-5'>自定義</b>軌跡規劃

    NVIDIA NeMo加速簡化自定義模型開發

    如果企業希望充分發揮出 AI 的力量,就需要根據其行業需求量身定制的自定義模型。
    的頭像 發表于 07-26 11:17 ?743次閱讀
    NVIDIA NeMo加速<b class='flag-5'>并</b>簡化<b class='flag-5'>自定義</b>模型開發

    如何手搓一個自定義的RPC 遠程過程調用框架

    是一種常用的技術,能夠簡化客戶端與服務器之間的交互。本文將介紹如何基于Netty(網絡編程框架)實現一個自定義的簡單的RPC框架。 首先簡單介紹一下RPC 主要特點: 1.1、RPC遠程過程調用的主要特點 ?透明性:
    的頭像 發表于 07-22 12:17 ?883次閱讀
    如何手搓一個<b class='flag-5'>自定義</b>的RPC 遠程過程<b class='flag-5'>調用</b>框架

    Labview實現自定義四維云圖(三維曲面圖像)可視化顯示

    函數基礎上,對其中的數據與程序進行修改,實現了曲面云圖顏色效果的自定義配置,提升了Labview的可視化應用性。本文以Labview自帶的“繪圖幫助(向量)”vi作為對象,詳細說明具體配置方法
    發表于 07-16 10:52 ?54次下載

    HarmonyOS開發案例:【 自定義彈窗】

    基于ArkTS的聲明式開發范式實現了三種不同的彈窗,第一種直接使用公共組件,后兩種使用CustomDialogController實現自定義彈窗
    的頭像 發表于 05-16 18:18 ?1353次閱讀
    HarmonyOS開發案例:【 <b class='flag-5'>自定義</b>彈窗】

    TSMaster 自定義 LIN 調度表編程指導

    LIN(LocalInterconnectNetwork)協議調度表是用于LIN總線通信中的消息調度的一種機制,我們收到越來越多來自不同用戶希望能夠通過接口實現自定義LIN調度表的需求。所以在
    的頭像 發表于 05-11 08:21 ?658次閱讀
    TSMaster <b class='flag-5'>自定義</b> LIN 調度表編程指導

    HarmonyOS實戰開發-深度探索與打造個性化自定義組件

    和可維護性,提高代碼效率。 自定義組件:是由@Component裝飾的UI單元,可以組合多個系統組件實現UI的復用,可以調用組件的生命周期。 自定義組件和頁面的關系 1.
    發表于 05-08 16:30

    HarmonyOS開發實例:【自定義Emitter】

    使用[Emitter]實現事件的訂閱和發布,使用[自定義彈窗]設置廣告信息。
    的頭像 發表于 04-14 11:37 ?995次閱讀
    HarmonyOS開發實例:【<b class='flag-5'>自定義</b>Emitter】

    Keil使用AC6編譯提示CUDA版本過高怎么解決?

    \' ArmClang: warning: Unknown CUDA version 10.2. Assuming the latest supported version 10.1
    發表于 04-11 07:56

    鴻蒙ArkUI實例:【自定義組件】

    組件是 OpenHarmony 頁面最小顯示單元,一個頁面可由多個組件組合而成,也可只由一個組件組合而成,這些組件可以是ArkUI開發框架自帶系統組件,比如?`Text`?、?`Button`?等,也可以是自定義組件,本節筆者簡單介紹一下自定義組件的語法規范。
    的頭像 發表于 04-08 10:17 ?632次閱讀

    深入淺出理解PagedAttention CUDA實現

    vLLM 中,LLM 推理的 prefill 階段 attention 計算使用第三方庫 xformers 的優化實現,decoding 階段 attention 計算則使用項目編譯 CUDA 代碼實現
    的頭像 發表于 01-09 11:43 ?1872次閱讀
    深入淺出理解PagedAttention <b class='flag-5'>CUDA</b><b class='flag-5'>實現</b>

    什么是CUDA?誰能打破CUDA的護城河?

    在最近的一場“AI Everywhere”發布會上,Intel的CEO Pat Gelsinger炮轟Nvidia的CUDA生態護城河并不深,而且已經成為行業的眾矢之的。
    的頭像 發表于 12-28 10:26 ?1.3w次閱讀
    什么是<b class='flag-5'>CUDA</b>?誰能打破<b class='flag-5'>CUDA</b>的護城河?

    基于YOLOv8實現自定義姿態評估模型訓練

    Hello大家好,今天給大家分享一下如何基于YOLOv8姿態評估模型,實現自定義數據集上,完成自定義姿態評估模型的訓練與推理。
    的頭像 發表于 12-25 11:29 ?2837次閱讀
    基于YOLOv8<b class='flag-5'>實現</b><b class='flag-5'>自定義</b>姿態評估模型訓練
    主站蜘蛛池模板: 免费a毛片| 视频一区亚洲中文字幕| 亚州AV中文无码乱人伦在线| 51成人精品午夜福利AV免费七| 国产精品久久久久久日本| 美女露出撒尿的部位| 亚洲成人网导航| 丰满五十老女人性视频| 老鸭窝毛片| 亚洲免费综合色视频| 纯肉小黄文高H| 末成年美女黄网站色大片连接| 亚洲av欧美在我| 国产AV亚洲精品久久久久| 免费成人高清在线视频| 亚洲欧美无码2017在线| 高清无码中文字幕影片| 欧美00后rapper潮水| 中国人泡妞xxxxxxxx19| 国产亚洲精品久久777777| 色悠久久综合| 草699一码二码三码四码| 蜜桃视频一区二区| 中文无码不卡的岛国片国产片| 国产色精品久久人妻无码看片软件| 青青草 久久久| 97国产视频| 李丽莎与土豪50分钟在线观看| 亚洲欧美在无码片一区二区 | 国产噜噜噜精品免费| 日本aa大片| avtt天堂网Av无码| 美女诱惑性感揉胸| 中文亚洲大香伊蕉不卡一区| 精品亚洲国产熟女福利自在线| 午夜免费体验30分| 国产第一页在线视频| 神马影院在线eecss伦理片| 大迪克黑人异族| 日韩一区二区在线免费观看| 把极品白丝班长啪到腿软|