提供統一記憶體架構:NVIDIA CUDA 6


話說,NVIDIA CUDA 這個 GPGPU 的程式開發架構從 2007 年推出 1.0 版發展至今,也已經好一段時間了;這段期間,NVIDIA 也不斷地推出新的 GPU 架構、以及對應的新版 CUDA SDK,在效能和功能上做強化,基本上應該也算是目前最成熟的 GPGPU 開發環境之一。

而日前,NVIDIA 也發表了還沒正式發布的最新版的 CUDA 6.0(現在還沒有可以下載的 SDK)的一個主要的功能,那就是「統一記憶體」(Unified Memory);原文可以參考官方的《Unified Memory in CUDA 6》一文。

在以往的 GPGPU 程式開發,不管怎樣,基本上都還是需要先把資料從 CPU 管理的系統記憶體(host、下圖左邊的 System Memory)、複製到顯示卡管理的記憶體(device、下圖右邊的 GPU Memory)上,然後才能開始使用 GPU 做計算;而計算完成之後,也需要手動把資料複製回 CPU、才能進一步處理這筆計算好的結果。

這個流程在硬體的架構上,基本上暫時應該還是無法避免的(AMD APU 的 HSA 應該算是特例);而 CUDA 6 的 Unified Memory,就是希望可以在程式開發上,可以針對這點做簡化、降低開發程式的門檻。

下圖就是 NVIDIA 提供的示意圖,左邊就是目前的狀況,而右邊則是引入 Unified Memory 的技術後的情況。

Unified Memory  基本上就是提供一個整合了 CPU 和 GPU 的記憶體管理區域,這裡面的記憶體可以以指標的形式、被 CPU 或 GPU 來做存取;系統會自動地在 host 與 device 之間,做記憶體搬移的動作,開發者不用再去自行做記憶體的複製動作。

這樣的好處,就是大幅地減低了記憶體管理上的複雜度,因為只要把這部分的工作都丟給系統去做就好了!開發者要做的,只是把配置記憶體的方法,從本來的 malloc() 換成 cudaMallocManaged(),釋放記憶體從 free() 變成 cudaFree() 就可以了。

下面就是官方提供的一個簡單的例子:

可以看到,在右方使用 CUDA 6 Unified Memory 的情況下,使用 cudaMallocManaged() 配置出來的記憶體,不但可以被 CUDA 的 kernel 執行,也可以被 CPU 上的 fread() 寫入資料。

而如果是要在 C++ 的環境下使用的話,基本上也可以透過自定義一個類別的建構子和解構子來做到更方便的管理;這部分可以參考《Unified Memory in CUDA 6》內的「Unified Memory with C++」一節,在這邊就不贅述了。

不過,這樣的自動記憶體管理,理論上應該還是會帶來一定程度的效能損失的。NVIDIA 到目前為止,都還沒有提及到這部分,所以使用這個技術到底對效能會有多大的影響,應該就要等 SDK 正式發布後、有人去試過才知道了;而到時候,是否要接受這樣的便利性/效能損失,就要看開發者的取捨了。

最後附帶一提,Unified Memory 只能用在 Kepler 架構的 GPU 上(Compute Capability 3.0+),使用的作業系統也有一些限制(64bit Windows 7+、Linux Kernel 2.6.18+),所以還在舊的開發環境的話…要用這項新功能的話,除了要等新的 SDK 發布、應該也等更新開發環境了~


除了 Unified Memory 外,CUDA 6 另外兩個被 highlight 出來的功能,還有:

  • Drop-in Libraries

    Automatically accelerate applications’ BLAS and FFTW calculations by up to 8X by simply replacing the existing CPU libraries with the GPU-accelerated equivalents.

  • Multi-GPU scaling

    The re-designed BLAS GPU library automatically scales performance across up to eight GPUs in a single node, delivering over nine teraflops of double precision performance per node, and supporting larger workloads than ever before (up to 512GB). The re-designed FFT GPU library scales up to 2 GPUs in a single node, allowing larger transform sizes and higher throughput.


nVidia CUDA 相關文章目錄

廣告

對「提供統一記憶體架構:NVIDIA CUDA 6」的想法

  1. 你好 這邊剛好拿到一張GTX760 測試了一下
    關於Cuda的統一記憶體架構測試後發現 基本上 data夠大還是會快
    但是 使用memcpy 跟 使用for迴圈來get access data 基本上慢很多
    目前 只是方便性質居多 要搞效能 可能還是有點悲劇
    vector1D*vector1D (size=10000000)
    這邊不是以你的計算方式 是比cpu速度
    統一架構
    以for迴圈 assign data( for(){cuda_a[i]=a[i]) )
    randdata=1215.881592

    Max error: 0 Average error: 0
    CPU : 137.885605 ms
    GPU(kernel) : 56.903584 ms (2.42x)
    GPU(kernel+mem) : 286.081848 ms (0.48x)

    memcpy device to device
    Max error: 0 Average error: 0
    CPU : 137.992294 ms
    GPU(kernel) : 1.229920 ms (112.20x)
    GPU(kernel+mem) : 36.661377 ms (3.76x)

    無統一架構memcpy host to device & device to host
    Randdata:1083.780273ms

    Max error: 0 Average error: 0
    OK!
    CPU : 80.397758 ms
    GPU : 1.189184 ms (67.61x)
    GPU(mem) : 39.109921 ms (2.06x)

    在產生data access 時的overhead好像也不小
    rand在統一架構下 rand慢了132ms
    cpu計算也受影響
    且假設如果需要在cpu使用data 可能還會再慢132ms?(猜測 沒真的算)

    結論大概就是直接access data 會比整塊存到device memory再回host 來的慢

    Liked by 1 person

    • Heresy 已經很久沒寫 CUDA 的程式了,UM 的架構也沒真的用過。
      不過,基本上,他就是為了便利性而出發的一個方案。

      但是,雖然在簡單的測試下,他的效能似乎比較差,不過在特定的狀況下,他似乎也可以有效能的增進。

      像在《An Investigation of Unified Memory Access Performance in CUDA》這篇裡面,就有提到:

      These results demonstrate that the performance of UMA varies greatly based on program design, kernel intensity, and memory access patterns. In particular, we demonstrate that there are indeed cases where a UMA version can perform better than a non-UMA implementation.

      http://www.bu.edu/peaclab/files/2014/03/TZhang_HPEC2014_UMA_camready.pdf

      而在《Unified Memory in CUDA 6: A Brief Overview and Related Data Access/Transfer Issues》這篇裡面,也可以看到部分的測試,UM 對效能是有增進的。
      http://sbel.wisc.edu/documents/TR-2014-09.pdf

      所以,應該還是看怎麼用、合不合適拿來用了。

  2. Hi, Heresy
    我目前在實現CUDA的運算,請問顯示卡的driver跟cuda memory的使用有絕對關係嗎??
    我嘗試使用簡單的小程式 – 將FHD的影像存到CPU memory 再寫到GPU的memory,然後再將GPU的memory資料寫CPU。但發現CPU的資料無法寫到GPU,用cudaError去檢查會發現錯誤,因為程式確認沒有問題,但由於某些原因不能更新driver的版本,請問您曾經遇過類似的問題嗎??

    謝謝~

    • CUDA SDK 應該是要對應驅動程式的版本的,如果是舊版驅動程式,應該無法正確對應新版的 CUDA SDK。
      如果你無法更新驅動程式的話,你可能要看看該版驅動程式能對應到哪個版本的 CUDA SDK 了。

發表迴響

在下方填入你的資料或按右方圖示以社群網站登入:

WordPress.com 標誌

您的留言將使用 WordPress.com 帳號。 登出 /  變更 )

Google+ photo

您的留言將使用 Google+ 帳號。 登出 /  變更 )

Twitter picture

您的留言將使用 Twitter 帳號。 登出 /  變更 )

Facebook照片

您的留言將使用 Facebook 帳號。 登出 /  變更 )

w

連結到 %s

This site uses Akismet to reduce spam. Learn how your comment data is processed.