在线观看www成人影院-在线观看www日本免费网站-在线观看www视频-在线观看操-欧美18在线-欧美1级

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

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

3天內不再提示

如何使用Warp在Python環境中編寫CUDA內核

星星科技指導員 ? 來源:NVIDIA ? 作者:NVIDIA ? 2022-04-02 16:15 ? 次閱讀

通常,實時物理模擬代碼是用低級 CUDA C ++編寫的,以獲得最佳性能。在這篇文章中,我們將介紹 NVIDIA Warp ,這是一個新的 Python 框架,可以輕松地用 Python 編寫可微圖形和模擬 GPU 代碼。 Warp 提供了編寫高性能仿真代碼所需的構建塊,但它的工作效率與 Python 等解釋語言相當。

在這篇文章的最后,您將學習如何使用 Warp 在 Python 環境中編寫 CUDA 內核,并利用一些內置的高級功能,從而輕松編寫復雜的物理模擬,例如海洋模擬。

安裝

Warp 以 來自 GitHub 的開源庫 的形式提供。克隆存儲庫后,可以使用本地軟件包管理器進行安裝。對于 pip ,請使用以下命令:

pip install warp

初始化

導入后,必須顯式初始化扭曲:

import warp as wp
wp.init()

推出內核

Warp 使用 Python 裝飾器的概念來標記可以在 GPU 上執行的函數。例如,可以編寫一個簡單的半隱式粒子積分方案,如下所示:

@wp.kernel
def integrate(x: wp.array(dtype=wp.vec3), v: wp.array(dtype=wp.vec3), f: wp.array(dtype=wp.vec3), w: wp.array(dtype=float), gravity: wp.vec3, dt: float): # thread id tid = wp.tid() x0 = x[tid] v0 = v[tid] # Semi-implicit Euler step f_ext = f[tid] inv_mass = w[tid] v1 = v0 + (f_ext * inv_mass + gravity) * dt x1 = x0 + v1 * dt # store results x[tid] = x1 v[tid] = v1 

因為 Warp 是強類型的,所以應該為內核參數提供類型提示。要啟動內核,請使用以下語法:

 wp.launch(kernel=simple_kernel, # kernel to launch dim=1024, # number of threads inputs=[a, b, c], # parameters device="cuda") # execution device

與 NumPy 等基于張量的框架不同, Warp 使用 kernel-based 編程模型。基于內核的編程與底層 GPU 執行模型更為匹配。對于需要細粒度條件邏輯和內存操作的模擬代碼,這通常是一種更自然的表達方式。然而, Warp 以一種易于使用的方式公開了這種以線程為中心的編程模型,它不需要 GPU 體系結構的低級知識。

編譯模型

啟動內核會觸發實時( JIT )編譯管道,該管道會自動從 Python 函數定義生成 C ++/ CUDA 內核代碼。

屬于 Python 模塊的所有內核都在運行時編譯到動態庫和 PTX 中。圖 2 。顯示了編譯管道,其中包括遍歷函數 AST 并將其轉換為直線 CUDA 代碼,然后編譯并加載回 Python 進程。

A flowchart diagram showing how Python code gets compiled and converted by Warp into kernel level executable code.A flowchart diagram showing how Python code gets compiled and converted by Warp into kernel level executable code.圖 2 。 Warp 內核的編譯管道

這個 JIT 編譯的結果被緩存。如果輸入內核源代碼不變,那么預編譯的二進制文件將以低開銷的方式加載。

記憶模型

Warp 中的內存分配通過warp.array類型公開。陣列封裝了可能位于主機( CPU )或設備( GPU )內存中的底層內存分配。與張量框架不同, Warp 中的數組是強類型的,并存儲內置結構的線性序列(vec3, matrix33, quat,等等)。

您可以從 Python 列表或 NumPy 數組中構造數組,或使用與 NumPy 和 PyTorch 類似的語法進行初始化:

# allocate an uninitizalized array of vec3s v = wp.empty(length=n, dtype=wp.vec3, device="cuda") # allocate a zero-initialized array of quaternions q = wp.zeros(length=n, dtype=wp.quat, device="cuda") # allocate and initialize an array from a numpy array # will be automatically transferred to the specified device v = wp.from_numpy(array, dtype=wp.vec3, device="cuda")

Warp 支持__array_interface____cuda_array_interface__協議,允許在基于張量的框架之間進行零拷貝數據視圖。例如,要將數據轉換為 NumPy ,請使用以下命令:

# automatically bring data from device back to host view = device_array.numpy()

特征

Warp 包含幾個更高級別的數據結構,使實現模擬和幾何處理算法更容易。

網格

三角形網格在仿真和計算機圖形學中無處不在。 Warp 提供了一種內置類型,用于管理網格數據,該數據支持幾何查詢,例如最近點、光線投射和重疊檢查。

下面的示例演示如何使用“扭曲”計算網格上距離輸入位置數組最近的點。這種類型的計算是碰撞檢測中許多算法的基礎(圖 3 )。 Warp 的網格查詢使實現此類方法變得簡單。

@wp.kernel
def project(positions: wp.array(dtype=wp.vec3), mesh: wp.uint64, output_pos: wp.array(dtype=wp.vec3), output_face: wp.array(dtype=int)): tid = wp.tid() x = wp.load(positions, tid) face_index = int(0) face_u = float(0.0) face_v = float(0.0) sign = float(0.0) max_dist = 2.0 if (wp.mesh_query_point(mesh, x, max_dist, sign, face_index, face_u, face_v)): p = wp.mesh_eval_position(mesh, face_index, face_u, face_v) output_pos[tid] = p output_face[tid] = face_index

稀疏卷

稀疏體對于表示大型域上的網格數據非常有用,例如復雜對象的符號距離場( SDF )或大規模流體流動的速度。 Warp 支持使用 NanoVDB 標準定義的稀疏卷。使用標準 OpenVDB 工具(如 Blender 、 Houdini 或 Maya )構造卷,然后在 Warp 內核內部采樣。

您可以直接從磁盤或內存中的二進制網格文件創建卷,然后使用volumes API 對其進行采樣:

wp.volume_sample_world(vol, xyz, mode) # world space sample using interpolation mode
wp.volume_sample_local(vol, uvw, mode) # volume space sample using interpolation mode
wp.volume_lookup(vol, ijk) # direct voxel lookup
wp.volume_transform(vol, xyz) # map point from voxel space to world space
wp.volume_transform_inv(vol, xyz) # map point from world space to volume space

使用卷查詢,您可以以最小的內存開銷高效地碰撞復雜對象。

散列網格

許多基于粒子的模擬方法,如離散元法( DEM )或平滑粒子流體動力學( SPH ),都涉及到在空間鄰域上迭代以計算力的相互作用。哈希網格是一種成熟的數據結構,用于加速這些最近鄰查詢,特別適合 GPU 。

哈希網格由點集構成,如下所示:

哈希網格由點集構成,如下所示:

grid = wp.HashGrid(dim_x=128, dim_y=128, dim_z=128, device="cuda")
grid.build(points=p, radius=r)

創建散列網格后,可以直接從用戶內核代碼中查詢它們,如以下示例所示,該示例計算所有相鄰粒子位置的總和:

@wp.kernel
def sum(grid : wp.uint64, points: wp.array(dtype=wp.vec3), output: wp.array(dtype=wp.vec3), radius: float): tid = wp.tid() # query point p = points[tid] # create grid query around point query = wp.hash_grid_query(grid, p, radius) index = int(0) sum = wp.vec3() while(wp.hash_grid_query_next(query, index)): neighbor = points[index] # compute distance to neighbor point dist = wp.length(p-neighbor) if (dist <= radius): sum += neighbor output[tid] = sum

圖 5 顯示了粘性材料的 DEM 顆粒材料模擬示例。使用內置的哈希網格數據結構,您可以在不到 200 行 Python 中編寫這樣的模擬,并以交互速率運行超過 100K 個粒子。

使用扭曲散列網格數據可以輕松評估相鄰粒子之間的成對力相互作用。

可微性

基于張量的框架,如 PyTorch 和 JAX ,提供了張量計算的梯度,非常適合于 ML 訓練等應用。

Warp 的一個獨特功能是能夠生成 kernel code 的正向和反向版本。這使得編寫可微模擬變得很容易,可以將梯度作為更大訓練管道的一部分進行傳播。一個常見的場景是,對網絡層使用傳統的 ML 框架,并使用 Warp 實現允許端到端差異性的模擬層。

當需要漸變時,應使用requires_grad=True創建陣列。例如,warp.Tape類可以記錄內核啟動并回放它們,以計算標量損失函數相對于內核輸入的梯度:

tape = wp.Tape() # forward pass
with tape: wp.launch(kernel=compute1, inputs=[a, b], device="cuda") wp.launch(kernel=compute2, inputs=[c, d], device="cuda") wp.launch(kernel=loss, inputs=[d, l], device="cuda") # reverse pass
tape.backward(loss=l)

完成后向傳遞后,可通過Tape對象中的映射獲得與輸入相關的梯度:

# gradient of loss with respect to input a
print(tape.gradients[a])
A 3D image with multi-colored trace lines simulating a ball bouncing off a wall and hitting a black square suspended in mid-air away from the wall.A 3D image with multi-colored trace lines simulating a ball bouncing off a wall and hitting a black square suspended in mid-air away from the wall.
圖 6 。一個軌跡優化的例子,其中球的初始速度被優化以擊中黑色目標。每行顯示 LBFGS 優化步驟的一次迭代的結果。

總結

在這篇文章中,我們介紹了 NVIDIA Warp ,這是一個 Python 框架,可以很容易地為 GPU 編寫可微模擬代碼。

關于作者

邁爾斯·麥克林( Miles Macklin )是NVIDIA 的首席工程師,致力于模擬技術。他從哥本哈根大學獲得計算機科學博士學位,從事計算機圖形學、基于物理學的動畫和機器人學的研究。他在 ACM SIGGRAPH 期刊上發表了幾篇論文,他的研究已經被整合到許多商業產品中,包括NVIDIA 的 PhysX 和 ISAAC 健身房模擬器。他最近的工作旨在為 GPU 上的可微編程開發健壯高效的框架。

Fred Oh 是 CUDA 、 CUDA on WSL 和 CUDA Python 的高級產品營銷經理。弗雷德擁有加州大學戴維斯分校計算機科學和數學學士學位。他的職業生涯開始于一名 UNIX 軟件工程師,負責將內核服務和設備驅動程序移植到 x86 體系結構。他喜歡《星球大戰》、《星際迷航》和 NBA 勇士隊。

審核編輯:郭婷

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

    關注

    211

    文章

    28419

    瀏覽量

    207108
  • NVIDIA
    +關注

    關注

    14

    文章

    4986

    瀏覽量

    103067
  • 自動駕駛
    +關注

    關注

    784

    文章

    13812

    瀏覽量

    166461
收藏 人收藏

    評論

    相關推薦

    晶圓的TTV,BOW,WARP,TIR是什么?

    晶圓的TTV、BOW、WARP、TIR是評估晶圓質量和加工精度的重要指標,以下是它們的詳細介紹: TTV(Total Thickness Variation,總厚度偏差) 定義:晶圓的總厚度變化
    的頭像 發表于 12-17 10:01 ?248次閱讀
    晶圓的TTV,BOW,<b class='flag-5'>WARP</b>,TIR是什么?

    邏輯異或運算符Python的用法

    Python的 ^ 符號實際上是一個按位異或運算符,用于對整數的二進制表示進行異或操作。 盡管如此,我們仍然可以通過一些方法來實現邏輯異或的功能,即當兩個布爾值不同時為真,相同時為假。這可以通過使用邏輯運算符來實現,而不是直接使用 ^ (因為 ^
    的頭像 發表于 11-19 09:46 ?186次閱讀

    Python環境下的代理服務器搭建與自動化管理

    Python環境下搭建與自動化管理代理服務器是一項涉及網絡編程和自動化技術的綜合任務。
    的頭像 發表于 11-14 07:31 ?168次閱讀

    怎么TMDSEVM6678: 6678自帶的FFT接口和CUDA提供CUFFT函數庫選擇?

    請教一下gpgpu上包括4個Riscv cpu和一個DPU, 沒有6678,要替換原來信號處理用的6678,該怎么6678自帶的FFT接口和CUDA提供CUFFT函數庫選擇?
    發表于 09-27 07:20

    linux驅動程序如何加載進內核

    Linux系統,驅動程序是內核與硬件設備之間的橋梁。它們允許內核與硬件設備進行通信,從而實現對硬件設備的控制和管理。 驅動程序的編寫
    的頭像 發表于 08-30 15:02 ?474次閱讀

    pytorch怎么pycharm運行

    PyTorch。以下是安裝PyTorch的步驟: 打開終端或命令提示符。 根據你的系統和需求,選擇適當的安裝命令。例如,如果你使用的是Python 3.8和CUDA 10.2,可以使用以下命令: pip
    的頭像 發表于 08-01 16:22 ?1426次閱讀

    PythonAI的應用實例

    Python人工智能(AI)領域的應用極為廣泛且深入,從基礎的數據處理、模型訓練到高級的應用部署,Python都扮演著至關重要的角色。以下將詳細探討Python
    的頭像 發表于 07-19 17:16 ?1098次閱讀

    打破英偉達CUDA壁壘?AMD顯卡現在也能無縫適配CUDA

    電子發燒友網報道(文/梁浩斌)一直以來,圍繞CUDA打造的軟件生態,是英偉達GPU領域最大的護城河,尤其是隨著目前AI領域的發展加速,市場火爆,英偉達GPU+CUDA的開發生態則更加穩固,AMD
    的頭像 發表于 07-19 00:16 ?4695次閱讀

    DongshanPI-AICT全志V853開發板搭建YOLOV5-V6.0環境

    的pytorch是否可用 Conda終端輸入python后,加載torch模塊,打印cuda是否可用。 (py37_yolov5) D:\\\\Programmers
    發表于 07-12 09:59

    用離線安裝器安裝的idf,其創建的Python虛擬環境無激活腳本是怎么回事?

    激活腳本,如何激活Python虛擬環境? 使用場景:Clion中用自定義腳本設置idf環境(Clion官方教程),附圖和鏈接如下: https://www.jetbrains.com
    發表于 06-11 06:49

    Win10 vscode無法編譯,提示python.exe: command not found怎么解決?

    電腦中已經刪除之前安裝的python環境變量也添加了報錯內容要求的python路徑,我用git bash、cmd也是可以打開python
    發表于 06-07 06:42

    AOSP源碼定制-內核驅動編寫

    有時候為了分析一些殼的檢測,需要在內核層面對讀寫相關的操作進行監控,每次去修改對應的內核源碼編譯重刷過于耗時耗力,這里就來嘗試編寫一個內核驅動,載入后監控讀寫。
    的頭像 發表于 04-23 11:15 ?1233次閱讀
    AOSP源碼定制-<b class='flag-5'>內核</b>驅動<b class='flag-5'>編寫</b>

    使用 PREEMPT_RT Ubuntu 構建實時 Linux 內核

    的實時內核補丁來完成。簡介我們曾介紹過Ubuntu22.04啟用實時Linux內核有多簡單,因為Canonical已將該內核列為一個選項
    的頭像 發表于 04-12 08:36 ?2473次閱讀
    使用 PREEMPT_RT <b class='flag-5'>在</b> Ubuntu <b class='flag-5'>中</b>構建實時 Linux <b class='flag-5'>內核</b>

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

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

    深入淺出理解PagedAttention CUDA實現

    vLLM ,LLM 推理的 prefill 階段 attention 計算使用第三方庫 xformers 的優化實現,decoding 階段 attention 計算則使用項目編譯 CUDA 代碼實現。
    的頭像 發表于 01-09 11:43 ?1893次閱讀
    深入淺出理解PagedAttention <b class='flag-5'>CUDA</b>實現
    主站蜘蛛池模板: 女人的逼毛片| 国产精品超清大白屁股 | 8x8x极品国产在线| 制服丝袜国产精品| 国产亚洲第一| 欧美精品 在线播放| 狠狠色狠狠色综合日日32| 深夜免费视频| 我被黑人巨大开嫩苞在线观看| 高清性欧美xxx| 中文字幕导航| 人人射人人| 天天做天天做天天综合网| 孩交精品xxxx视频视频| 免费观看激色视频网站bd| 天天碰免费视频| 最新欧美一级视频| 91成人免费| 欧美性狂猛bbbbbxxxxx| 成人小视频在线| 三级成人网| 在线资源天堂| 日本一线a视频免费观看| 好男人社区www的视频免费| 激情五月激情综合网| 伊人精品网| 嫩草影院在线入口| 久久天天躁狠狠躁夜夜2020一| 一区二区视频| 欧美精品1| 福利视频网址| 三级黄色网址| 最新福利网站| 日本一区二区在线视频| 三级在线看| 色综合视频| 亚洲国产成人精品青青草原100| bt天堂在线www最新版资源网| 国产大片免费观看资源| 四虎成人精品在永久在线观看| 亚洲免费色视频|