CUDA 學習筆記 (一) 【Chapter3 CUDA初探】
Chapter3
目錄
3.1 hello , world !
/* hello_world.cu*/
#include "../common/book.h"
int main( void ) {
printf( "Hello, World!\n" );
return 0;
}這個簡單的 hello world程式只是為了說明， CUDA C與你熟悉的標準C在很大程度上是沒有區別的。這個示例很簡單，它能夠完全在主機上執行。這個示例引出了本書的一個重要區分：我們將CPU以及系統的記憶體
接下來將逐漸完善這個簡單示例。我們來看看如何使用GPU（這就是一個裝置）來執行程式碼。在GPU裝置上執行的函式通常稱為核函式（ Kernel）。
3.2 核函式呼叫
/* simple_kernel.cu */ #include "../common/book.h" __global__ void kernel( void ) { } int main( void ) { kernel<<<1,1>>>(); printf( "Hello, World!\n" ); return 0; }
這個程式與最初的“ Hello， World！”相比，多了兩個值得注意的地方：
- 一個空的函式 kernel()，並且帶有修飾符 __global__。
- 對這個空函式的呼叫，並且帶有修飾字符 <<<1,1>>>。
CUDA C為標準C增加了__global__修飾符，__global__修飾符將告訴編譯器kernel()函式應該編譯為在裝置上執行。在這個簡單的示例中，函式kernel()將被交給編譯裝置程式碼的編譯器，而main()函式將被交給主機編譯器（與上一個例子一樣）。
那麼， kernel()的呼叫究竟代表著什麼含義，並且為什麼必須加上尖括號和兩個數值？注意，這正是使用CUDA C的地方。
我們已經看到， CUDA C需要通過某種語法方法將一個函式標記為“裝置程式碼（ DeviceCode）”。這並沒有什麼特別之處，而只是一種簡單的表示方法，表示將主機程式碼傳送到一個編譯器，而將裝置程式碼傳送到另一個編譯器。事實上，這裡的關鍵在於如何在主機程式碼中呼叫裝置程式碼。 CUDA C的優勢之一在於，它提供了與C在語言級別上的整合，因此這個裝置函式呼叫看上去非常像主機函式呼叫。在後面將詳細介紹在這個函式呼叫背後發生的動作，但就目前而言，只需知道CUDA編譯器和執行時將負責實現從主機程式碼中呼叫裝置程式碼。
因此，這個看上去有些奇怪的函式呼叫實際上表示呼叫裝置程式碼，但為什麼要使用尖括號和數字？尖括號表示要將一些引數傳遞給執行時系統。這些引數並不是傳遞給裝置程式碼的引數，而是告訴執行時如何啟動裝置程式碼。在第4章中，我們將瞭解這些引數對執行時的作用。傳遞給裝置程式碼本身的引數是放在圓括號中傳遞的，就像標準的函式呼叫一樣。
3.3 傳遞引數
將引數傳遞給核函式的簡單示例：
/* simple_device_call.cu */
#include "../common/book.h"
__device__ int addem( int a, int b ) {
return a + b;
}
__global__ void add( int a, int b, int *c ) {
*c = addem( a, b );
}
int main( void ) {
int c;
int *dev_c;
HANDLE_ERROR( cudaMalloc( (void**)&dev_c, sizeof(int) ) );
add<<<1,1>>>( 2, 7, dev_c );
HANDLE_ERROR( cudaMemcpy( &c,
dev_c,
sizeof(int),
cudaMemcpyDeviceToHost ) );
printf( "2 + 7 = %d\n", c );
HANDLE_ERROR( cudaFree( dev_c ) );
return 0;
}注意這裡增加了多行程式碼，在這些程式碼中包含兩個概念：
- 可以像呼叫C函式那樣將引數傳遞給核函式。
- 當裝置執行任何有用的操作時，都需要分配記憶體，例如將計算值返回給主機。
將引數傳遞給核函式：傳參過程除了尖括號語法之外，核函式的外表和行為看上去與標準C中的任何函式呼叫一樣。執行時系統負責處理將引數從主機傳遞給裝置的過程中的所有複雜操作。
如何在裝置上分配記憶體和釋放記憶體：
- cudaMalloc()來分配裝置記憶體
- cudaFree()釋放裝置記憶體
- 注意：不能在主機程式碼中使用cudaMalloc()分配的指標進行記憶體讀/寫操作
通過cudaMalloc()來分配記憶體：類似於標準的C函式malloc()，cudaMalloc()函式的作用是告訴CUDA執行時在裝置上分配記憶體。cudaMalloc() 函式有兩個引數：第一個引數是一個指標（指向用於儲存新分配記憶體地址的變數）；第二個引數是分配記憶體的大小。除了分配記憶體的指標不是作為函式的返回值外，這個函式的行為與malloc()是相同的，並且返回型別為void*。
【注意】CUDA C的簡單性及其強大功能在很大程度上都是來源於它淡化了主機程式碼和裝置程式碼之間的差異。然而，程式設計師一定不能在主機程式碼中對cudaMalloc()返回的指標進行解引用（ Dereference）。主機程式碼可以將這個指標作為引數傳遞，對其執行算術運算，甚至可以將其轉換為另一種不同的型別。但是，絕對不可以使用這個指標來讀取或者寫入記憶體。遺憾的是，編譯器無法防止這種錯誤的發生。我們可以將裝置指標的使用限制總結如下：
- 可以將cudaMalloc()分配的指標傳遞給在裝置上執行的函式。
- 可以在裝置程式碼中使用cudaMalloc()分配的指標進行記憶體讀/寫操作。
- 可以將cudaMalloc()分配的指標傳遞給在主機上執行的函式。
- 不能在主機程式碼中使用cudaMalloc()分配的指標進行記憶體讀/寫操作。
- 不能使用標準C的 free() 函式來釋放cudaMalloc()分配的記憶體。需要呼叫cudaFree()釋放cudaMalloc()分配的記憶體（這個函式的行為與free()的行為非常相似）。
HANDLE_ERROR()：函式呼叫外層的HANDLE_ERROR()是我們定義的一個巨集，作為本書輔助程式碼的一部分。這個巨集只是判斷函式呼叫是否返回了一個錯誤值，如果是的話，那麼將輸出相應的錯誤訊息，退出應用程式並將退出碼設定為EXIT_FAILURE。雖然你也可以在自己的應用程式中使用這個錯誤處理碼，但這種做法在產品級的程式碼中很可能是不夠的。
訪問裝置記憶體----使用裝置指標 / 呼叫cudaMemcpy()函式：
- 在裝置程式碼中使用裝置指標訪問裝置記憶體
- 注意：雖然可以將主機指標傳遞給裝置程式碼，不能通過主機指標訪問裝置的記憶體
- 在主機程式碼中呼叫cudaMemcpy()來訪問裝置上的記憶體
在裝置程式碼中使用裝置指標：
裝置指標的使用方式與標準C中指標的使用方式完全一樣。語句*c = a + b的含義同樣非常簡單：將引數a和b相加，並將結果儲存在c指向的記憶體中。在前面列出了裝置指標的使用限制，主機指標的使用也有著類似的限制。雖然可以將主機指標傳遞給裝置程式碼，不能通過主機指標訪問裝置的記憶體。 總之，主機指標只能訪問主機程式碼中的記憶體，而裝置指標也只能訪問裝置程式碼中的記憶體。
呼叫cudaMemcpy()：在主機程式碼中可以通過呼叫cudaMemcpy()來訪問裝置上的記憶體。這個函式呼叫的行為類似於標準C中的memcpy()，只不過多了一個引數來指定裝置記憶體指標究竟是源指標還是目標指標。
【注意】cudaMemcpy()的引數：
- cudaMemcpyDeviceToHost 將告訴執行時源指標是一個裝置指標，而目標指標是一個主機指標（本示例即是如此）
- cudaMemcpyHostToDevice 將告訴執行時源指標是一個主機指標，而目標指標是一個裝置指標。
- cudaMemcpyDeviceToDevice 將告訴執行時兩個指標都是位於裝置上。
- 如果源指標和目標指標都位於主機上，那麼可以直接呼叫標準C的 memcpy() 函式。
3.4 查詢裝置
由於我們希望在裝置上分配記憶體和執行程式碼，因此如果在程式中能夠知道裝置擁有多少記憶體以及具備哪些功能，那麼將非常有用。而且，在一臺計算機上擁有多個支援CUDA的裝置也是很常見的情形。在這些情況中，我們希望通過某種方式來確定使用的是哪一個處理器。例如，在許多主機板中都集成了NVIDIA圖形處理器。當計算機生產商或者使用者將一塊獨立的圖形處理器新增到計算機時，那麼就有了兩個支援CUDA的處理器。某些NVIDIA產品，例如GeForce GTX 295，在單塊卡上包含了兩個GPU，因此使用這類產品的計算機也就擁有了兩個支援CUDA的處理器。
在深入研究如何編寫裝置程式碼之前，我們需要通過某種機制來判斷計算機中當前有哪些裝置，以及每個裝置都支援哪些功能。幸運的是，可以通過一個非常簡單的介面來獲得這種資訊。首先，我們希望知道在系統中有多少個裝置是支援CUDA架構的，並且這些裝置能夠執行基於CUDA C編寫的核函式。要獲得CUDA裝置的數量，可以呼叫cudaGetDeviceCount()。這個函式的作用從它的名字就可以看出來。
int count;
HANDLE_ERROR( cudaGetDeviceCount( &count ) );
在呼叫cudaGetDeviceCount()後，可以對每個裝置進行迭代，並查詢各個裝置的相關資訊。CUDA執行時將返回一個cudaDeviceProp型別的結構，其中包含了裝置的相關屬性。我們可以獲得哪些屬性？從CUDA 3.0開始，在cudaDeviceProp結構中包含了以下資訊：
struct cudaDeviceProp
{
char name[256];
size_t totalGlobalMem;
size_t sharedMemPerBlock;
int regsPerBlock;
int warpSize;
size_t memPitch;
int maxThreadsPerBlock;
int maxThreadsDim[3];
int maxGridSize[3];
size_t totalConstMem;
int major;
int minor;
int clockRate;
size_t textureAlignment;
int deviceOverlap;
int multiProcessorCount;
int kernelExecTimeoutEnabled;
int integrated;
int canMapHostMemory;
int computeMode;
int maxTexture1D;
int maxTexture2D[2];
int maxTexture3D[3];
int maxTexture2DArray[3];
int concurrentKernels;
}
其中，有些屬性的含義是顯而易見的，其他屬性的含義如下所示（見表3.1）。
表3.1 CUDA裝置屬性
--------------------------------------------------------------------------------------------------------------------------------------------------------------
設 備 屬 性 描 述
--------------------------------------------------------------------------------------------------------------------------------------------------------------
char name[256]; 標識裝置的ASCII字串（例如， "GeForce GTX 280"）
size_t totalGlobalMem 裝置上全域性記憶體的總量，單位為位元組
size_t sharedMemPerBlock 在一個執行緒塊（ Block）中可使用的最大共享記憶體數量，單位為位元組
int regsPerBlock 每個執行緒塊中可用的32位暫存器數量
int warpSize 在一個執行緒束（ Warp）中包含的執行緒數量
size_t memPitch 在記憶體複製中最大的修正量（ Pitch），單位為位元組
int maxThreadsPerBlock 在一個執行緒塊中可以包含的最大執行緒數量
int maxThreadsDim[3] 在多維執行緒塊陣列中，每一維可以包含的最大執行緒數量
int maxGridSize[3] 在一個執行緒格（ Grid）中，每一維可以包含的執行緒塊數量
size_t totalConstMem 常量記憶體的總量
int major 裝置計算功能集（ Compute Capability）的主版本號
int minor 裝置計算功能集的次版本號
size_t textureAlignment 裝置的紋理對齊（ Texture Alignment）要求
int deviceOverlap 一個布林型別值，表示裝置是否可以同時執行一個cudaMemory()呼叫和一個核函式呼叫
int multiProcessorCount 裝置上多處理器的數量
int kernelExecTimeoutEnabled 一個布林值，表示在該裝置上執行的核函式是否存在執行時限制
int integrated 一個布林值，表示裝置是否是一個整合GPU（即該GPU屬於晶片組的一部分而非獨立的GPU）
int canMapHostMemory 一個布林型別的值，表示裝置是否將主機記憶體對映到CUDA裝置地址空間
int computeMode 表示裝置的計算模式：預設（ Default），獨佔（ Exclusive），或者禁止（ Prohibited）
int maxTexture1D 一維紋理的最大大小
int maxTexture2D[2] 二維紋理的最大維數
int maxTexture3D[3] 三維紋理的最大維數
int maxTexture2DArray[3] 二維紋理陣列的最大維數
int concurrentKernels 一個布林型別值，表示裝置是否支援在同一個上下文中同時執行多個核函式
--------------------------------------------------------------------------------------------------------------------------------------------------------------
就目前而言，我們不會詳細介紹所有這些屬性。事實上，在上面的列表中沒有給出屬性的一些重要細節，因此你需要參考《 NVIDIA CUDA Programming Guide》以瞭解更多的資訊。當開始編寫應用程式時，這些屬性會非常有用。但就目前而言，我們只是給出瞭如何查詢每個裝置並且報告裝置的相應屬性。下面給出了對裝置進行查詢的程式碼：
/* searchDevice.cu */
#include "../common/book.h"
int main( void ) {
cudaDeviceProp prop;
int dev;
HANDLE_ERROR(cudaGetDevice( &dev ));
printf( "ID of current CUDA device: %d\n", dev );
memset( &prop, 0, sizeof( cudaDeviceProp ) );
prop.major = 1;
prop.minor = 3;
HANDLE_ERROR( cudaChooseDevice( &dev, &prop ) );
printf( "ID of CUDA device closest to revision 1.3: %d\n", dev );
HANDLE_ERROR( cudaSetDevice( dev ) );
}
在知道了每個可用的屬性後，接下來就可以將註釋“對裝置的屬性執行某些操作”替換為一些具體的操作：
/* enum_gpu.cu */
#include "../common/book.h"
int main( void ) {
cudaDeviceProp prop;
int count;
HANDLE_ERROR( cudaGetDeviceCount( &count ) );
for (int i=0; i< count; i++) {
HANDLE_ERROR( cudaGetDeviceProperties( &prop, i ) );
printf( " --- General Information for device %d ---\n", i );
printf( "Name: %s\n", prop.name );
printf( "Compute capability: %d.%d\n", prop.major, prop.minor );
printf( "Clock rate: %d\n", prop.clockRate );
printf( "Device copy overlap: " );
if (prop.deviceOverlap)
printf( "Enabled\n" );
else
printf( "Disabled\n");
printf( "Kernel execution timeout : " );
if (prop.kernelExecTimeoutEnabled)
printf( "Enabled\n" );
else
printf( "Disabled\n" );
printf( " --- Memory Information for device %d ---\n", i );
printf( "Total global mem: %ld\n", prop.totalGlobalMem );
printf( "Total constant Mem: %ld\n", prop.totalConstMem );
printf( "Max mem pitch: %ld\n", prop.memPitch );
printf( "Texture Alignment: %ld\n", prop.textureAlignment );
printf( " --- MP Information for device %d ---\n", i );
printf( "Multiprocessor count: %d\n",
prop.multiProcessorCount );
printf( "Shared mem per mp: %ld\n", prop.sharedMemPerBlock );
printf( "Registers per mp: %d\n", prop.regsPerBlock );
printf( "Threads in warp: %d\n", prop.warpSize );
printf( "Max threads per block: %d\n",
prop.maxThreadsPerBlock );
printf( "Max thread dimensions: (%d, %d, %d)\n",
prop.maxThreadsDim[0], prop.maxThreadsDim[1],
prop.maxThreadsDim[2] );
printf( "Max grid dimensions: (%d, %d, %d)\n",
prop.maxGridSize[0], prop.maxGridSize[1],
prop.maxGridSize[2] );
printf( "\n" );
}
}
3.5 裝置屬性的使用
除非是編寫一個需要輸出每個支援CUDA的顯示卡的詳細屬性的應用程式，否則我們是否需要了解系統中每個裝置的屬性？作為軟體開發人員，我們希望編寫出的軟體是最快的，因此可能需要選擇擁有最多處理器的GPU來執行程式碼。或者，如果核函式與CPU之間需要進行密集互動，那麼可能需要在整合的GPU上執行程式碼，因為它可以與CPU共享記憶體。這兩個屬性都可以通過cudaGetDeviceProperties()來查詢。
假設我們正在編寫一個需要使用雙精度浮點計算的應用程式。在快速翻閱《 NVIDIA CUDA Programming Guide》的附錄A後，我們知道計算功能集的版本為1.3或者更高的顯示卡才能支援雙精度浮點數學計算。因此，要想成功地在應用程式中執行雙精度浮點運算， GPU裝置至少需要支援1.3或者更高版本的計算功能集。
根據在 cudaGetDeviceCount()和 cudaGetDeviceProperties()中返回的結果，我們可以對每個裝置進行迭代，並且查詢主版本號大於1，或者主版本號為1且次版本號大於等於3的裝置。但是，這種迭代操作執行起來有些繁瑣，因此CUDA執行時提供了一種自動方式來執行這個迭代操作。首先，找出我們希望裝置擁有的屬性並將這些屬性填充到一個cudaDeviceProp結構。
cudaDeviceProp prop;
memset( &prop, 0, sizeof( cudaDeviceProp ) );
prop.major = 1;
prop.minor = 3;在填充完 cudaDeviceProp 結構後，將其傳遞給 cudaChooseDevice()，這樣CUDA執行時將查詢是否存在某個裝置滿足這些條件。 cudaChooseDevice()函式將返回一個裝置ID，然後我們可以將這個ID傳遞給 cudaSetDevice()。隨後，所有的裝置操作都將在這個裝置上執行。
/* set_gpu.cu */
#include "../common/book.h"
int main( void ) {
cudaDeviceProp prop;
int dev;
HANDLE_ERROR( cudaGetDevice( &dev ) );
printf( "ID of current CUDA device: %d\n", dev );
memset( &prop, 0, sizeof( cudaDeviceProp ) );
prop.major = 1;
prop.minor = 3;
HANDLE_ERROR( cudaChooseDevice( &dev, &prop ) );
printf( "ID of CUDA device closest to revision 1.3: %d\n", dev );
HANDLE_ERROR( cudaSetDevice( dev ) );
}
當前，在系統中擁有多個GPU已是很常見的情況。例如，許多NVIDIA主機板晶片組都包含了整合的並且支援CUDA的GPU。當把一個獨立的GPU新增到這些系統中時，那麼就形成了一個多GPU的平臺。而且， NVIDIA的SLI（Scalable Link Interface，可伸縮鏈路介面）技術使得多個獨立的GPU可以並排排列。無論是哪種情況，應用程式都可以從多個GPU中選擇最適合的GPU。如果應用程式依賴於GPU的某些特定屬性，或者需要在系統中最快的GPU上執行，那麼你就需要熟悉這個API，因為CUDA執行時本身並不能保證為應用程式選擇最優或者最合適的GPU。
3.6 本章小結
從本質上來說，CUDA C只是對標準C進行了語言級的擴充套件，通過增加一些修飾符使我們可以指定哪些程式碼在裝置上執行，以及哪些程式碼在主機上執行。在函式前面新增關鍵字__global__將告訴編譯器把該函式放在GPU上執行。為了使用GPU的專門記憶體，我們還學習了與C的 malloc()， memcpy()和 free()等API對應的CUDA API。這些函式的CUDA版本，包括cudaMalloc(), cudaMemcpy()以及 cudaFree()，分別實現了分配裝置記憶體，在裝置和主機之間複製資料，以及釋放裝置記憶體等功能。後面還將介紹一些更有趣的示例，這些示例都是關於如何將GPU裝置作為一種大規模並行協處理器來使用。
參考書：《GPU高效能程式設計CUDA實戰》
參考blog: https://blog.csdn.net/w09103419/article/details/52484969ca