【轉】cuda 《GPU高性能編程CUDA實戰》中代碼整理


    CUDA架構專門為GPU計算設計了一種全新的模塊,目的是減輕早期GPU計算中存在的一些限制,而正是這些限制使得之前的GPU在通用計算中沒有得到廣泛的應用。

         使用CUDA C來編寫代碼的前提條件包括:(1)、支持CUDA的圖形處理器,即由NVIDIA推出的GPU顯卡,要求顯存超過256MB;(2)、NVIDIA設備驅動程序,用於實現應用程序與支持CUDA的硬件之間的通信,確保安裝最新的驅動程序,注意選擇與開發環境相符的圖形卡和操作系統;(3)、CUDA開發工具箱即CUDA Toolkit,此工具箱中包括一個編譯GPU代碼的編譯器;(4)、標准C編譯器,即CPU編譯器。CUDA C應用程序將在兩個不同的處理器上執行計算,因此需要兩個編譯器。其中一個編譯器為GPU編譯代碼,而另一個為CPU編譯代碼。 

         一般,將CPU以及系統的內存稱為主機(Host),而將GPU及其內存稱為設備(Device)。在GPU設備上執行的函數通常稱為核函數(Kernel)

         cudaMalloc函數使用限制總結:(1)、可以將cudaMalloc()分配的指針傳遞給在設備上執行的函數;(2)、可以在設備代碼中使用cudaMalloc()分配的指針進行內存讀/寫操作;(3)、可以將cudaMalloc()分配的指針傳遞給在主機上執行的函數;(4)、不能在主機代碼中使用cudaMalloc()分配的指針進行內存讀/寫操作。

         不能使用標准C的free()函數來釋放cudaMalloc()分配的內存;要釋放cudaMalloc()分配的內存,需要調用cudaFree()。

         設備指針的使用方式與標准C中指針的使用方式完全一樣。主機指針只能訪問主機代碼中的內存,而設備指針也只能訪問設備代碼中的內存。

         在主機代碼中可以通過調用cudaMemcpy()來訪問設備上的內存。

         有可能在單塊卡上包含了兩個或多個GPU。

         在集成的GPU上運行代碼,可以與CPU共享內存

         計算功能集的版本為1.3或者更高的顯卡才能支持雙精度浮點數的計算

         尖括號的第一個參數表示設備在執行核函數時使用的並行線程塊的數量,

         並行線程塊集合也稱為一個線程格(Grid)。線程格既可以是一維的線程塊集合,也可以是二維的線程塊集合。

         GPU有着完善的內存管理機制,它將強行結束所有違反內存訪問規則的進程

         在啟動線程塊數組時,數組每一維的最大數量都不能超過65535.這是一種硬件限制,如果啟動的線程塊數量超過了這個限值,那么程序將運行失敗。   

         CUDA運行時將線程塊(Block)分解為多個線程。當需要啟動多個並行線程塊時,只需將尖括號中的第一個參數由1改為想要啟動的線程塊數量。在尖括號中,第二個參數表示CUDA運行時在每個線程塊中創建的線程數量。

         硬件將線程塊的數量限制為不超過65535.同樣,對於啟動核函數時每個線程塊中的線程數量,硬件也進行了限制。具體來說,最大的線程數量不能超過設備屬性結構中maxThreadsPerBlock域的值。這個值並不固定,有的是512,有的是1024.

         內置變量blockDim,對於所有線程塊來說,這個變量是一個常數,保存的是線程塊中每一維的線程數量。

         內置變量gridDim,對於所有線程塊來說,這個變量是一個常數,用來保存線程格每一維的大小,即每個線程格中線程塊的數量。

         內置變量blockIdx,變量中包含的值就是當前執行設備代碼的線程塊的索引。

         內置變量threadIdx,變量中包含的值就是當前執行設備代碼的線程索引。

         CUDA運行時允許啟動一個二維線程格,並且線程格中的每個線程塊都是一個三維的線程數組。

         CUDA C支持共享內存:可以將CUDA  C的關鍵字__share__添加到變量聲明中,這將使這個變量駐留在共享內存中。CUDA C編譯器對共享內存中的變量與普通變量將分別采取不同的處理方式。

         CUDA架構將確保,除非線程塊中的每個線程都執行了__syncthreads(),否則沒有任何線程能執行__syncthreads()之后的指令。

         由於在GPU上包含有數百個數學計算單元,因此性能瓶頸通常並不在於芯片的數學計算吞吐量,而是在於芯片的內存帶寬。

         常量內存用於保存在核函數執行期間不會發生變化的數據。NVIDIA硬件提供了64KB的常量內存,並且對常量內存采取了不同於標准全局內存的處理方式。在某些情況下,用常量內存來替換全局內存能有效地減少內存帶寬。要使用常量內存,需在變量前面加上__constant__關鍵字。

         在CUDA架構中,線程束是指一個包含32個線程的集合,這個線程集合被“編織在一起”並且以“步調一致(Lockstep)”的形式執行。在程序中的每一行,線程束中的每個線程都將在不同的數據上執行相同的指令。

         紋理內存是在CUDA C程序中可以使用的另一種只讀內存。與常量內存類似的是,紋理內存同樣緩存在芯片上,因此在某些情況中,它能夠減少對內存的請求並提供更高效的內存帶寬。紋理緩存是專門為那些在內存訪問模式中存在大量空間局部性(Spatial Locality)的圖形應用程序而設計的。

         NVIDIA將GPU支持的各種功能統稱為計算功能集(Compute Capability)。高版本計算功能集是低版本計算功能集的超集

         只有1.1或者更高版本的GPU計算功能集才能支持全局內存上的原子操作。此外,只有1.2或者更高版本的GPU計算功能集才能支持共享內存上的原子操作。CUDA C支持多種原子操作。

         C庫函數malloc函數將分配標准的,可分頁的(Pageble)主機內存。而cudaHostAlloc函數將分配頁鎖定的主機內存。頁鎖定內存也稱為固定內存(Pinned Memory)或者不可分頁內存,它有一個重要的屬性:操作系統將不會對這塊內存分頁並交換到磁盤上,從而確保了該內存始終駐留在物理內存上。因此,操作系統能夠安全地使某個應用程序訪問該內存的物理地址,因為這塊內存將不會被破壞或者重新定位。

         固定內存是一把雙刃劍。當使用固定內存時,你將失去虛擬內存的所有功能。特別是,在應用程序中使用每個頁鎖定內存時都需要分配物理內存,因為這些內存不能交換到磁盤上。這意味着,與使用標准的malloc函數調用相比,系統將更快地耗盡內存。因此,應用程序在物理內存較少的機器上會運行失敗,而且意味着應用程序將影響在系統上運行的其它應用程序的性能。

         建議,僅對cudaMemcpy()調用中的源內存或者目標內存,才使用頁鎖定內存,並且在不再需要使用它們時立即釋放,而不是等到應用程序關閉時才釋放。

         CUDA流表示一個GPU操作隊列,並且該隊列中的操作將以指定的順序執行

         通過使用零拷貝內存,可以避免CPU和GPU之間的顯式復制操作

         對於零拷貝內存,獨立GPU和集成GPU,帶來的性能提升是不同的。對於集成GPU,使用零拷貝內存通常都會帶來性能提升,因為內存在物理上與主機是共享的。將緩沖區聲明為零拷貝內存的唯一作用就是避免不必要的數據復制。但是,所有類型的固定內存都存在一定的局限性,零拷貝內存同樣不例外。每個固定內存都會占用系統的可用物理內存,這最終將降低系統的性能。對於獨立GPU,當輸入內存和輸出內存都只能使用一次時,那么在獨立GPU上使用零拷貝內存將帶來性能提升。但由於GPU不會緩存零拷貝內存的內容,如果多次讀取內存,那么最終將得不償失,還不如一開始就將數據復制到GPU。

         CUDA工具箱(CUDAToolkit)包含了兩個重要的工具庫:(1)、CUFFT(Fast FourierTransform,快速傅里葉變換)庫;(2)、CUBLAS(Basic Linear Algebra Subprograms,BLAS)是一個線性代數函數庫。

         NPP(NVIDIA Performance Primitives)稱為NVIDIA性能原語,它是一個函數庫,用來執行基於CUDA加速的數據處理操作,它的基本功能集合主要側重於圖像處理和視頻處理。

新建一個基於CUDA的測試工程testCUDA,此工程中除了包括common文件外,還添加了另外三個文件,分別為testCUDA.cu、funset.cu、funset.cuh,這三個文件包括了書中絕大部分的測試代碼:


免責聲明!

本站轉載的文章為個人學習借鑒使用,本站對版權不負任何法律責任。如果侵犯了您的隱私權益,請聯系本站郵箱yoyou2525@163.com刪除。



 
粵ICP備18138465號   © 2018-2025 CODEPRJ.COM