前言
CUDA並行程序設計系列是本人在學習CUDA時整理的資料，內容大都來源於對《CUDA並行程序設計：GPU編程指南》、《GPU高性能編程CUDA實戰》和CUDA Toolkit Documentation的整理。通過本系列整體介紹CUDA並行程序設計。內容包括GPU簡介、CUDA簡介、環境搭建、線程模型、內存、原子操作、同步、流和多GPU架構等。
本系列目錄：
- 【CUDA並行程序設計系列（1）】GPU技術簡介
- 【CUDA並行程序設計系列（2）】CUDA簡介及CUDA初步編程
- 【CUDA並行程序設計系列（3）】CUDA線程模型
- 【CUDA並行程序設計系列（4）】CUDA內存
- 【CUDA並行程序設計系列（5）】CUDA原子操作與同步
- 【CUDA並行程序設計系列（6）】CUDA流與多GPU
- 關於CUDA的一些學習資料
本文對CUDA進行簡單介紹，並通過實例代碼演示怎么編寫在GPU上運行的代碼，最后寫一段代碼來查詢本機GPU的設備參數。
CUDA C簡介
CUDA C是NVIDIA公司設計的一種編程語言，用於在GPU上得編寫通用計算程序，目前也叫“ CUDA C and C++”，或者“CUDA C/C++”。CUDA C是C語言的擴展，所以使用起來和C語言類似。當然，CUDA現在已經不局限於C語言了，在NVIDIA ZONE的LANGUAGE SOLUTIONS就明確支持多種語言和開發環境，如：C++、Python、Java、.Net、OpenACC、OpenCL等，操作系統也支持Linux、Mac OS、Windows。當然，前提是電腦至少配備一個支持CUDA的GPU，在NVIDIA官方可以查看自己電腦的顯卡是否支持CUDA。
GPU計算能力
在查看某個顯卡是否支持CUDA時，還會看到一個計算能力（Compute Capability）的參數。正如不同的CPU有着不同的功能和指令集，對於支持CUDA的GPU也同樣如此。NVIDIA將GPU支持的各種功能統稱為計算能力。硬件的計算能力是固定的。不同計算能力具有一定差別，如，在計算能力1.0版本不支持全局內存上的原子操作。更高計算能力的GPU是低計算能力的超級，所以計算能力2.0支持的功能在3.0也全部支持。可以看到，目前GPU最高的計算能力已經達到5.3（Tegra X1）。
環境搭建
環境搭建比較簡單，在NVIDIA開發官網選擇對應的平台下載最新版（目前是7.5）並安裝就行了。在CUDA Toolkit Online Documentation有詳細的安裝教程。
Hello,World!
新建一個文件“hello_world.cu"
int main( void )
{
printf( "Hello, World!\n" );
return 0;
}
除了后綴”.cu”表示CUDA文件，這段代碼甚至不需要任何解釋，下面編譯並運行：
nvcc hello_world.cu
./a.out
輸出：
Hello, World!
我們使用nvcc命令編譯，這將使用CUDA C編譯器來編譯這段代碼。
修改一下這段代碼：
__global__ void kernel( void )
{
}
int main( void )
{
kernel<<<1,1>>>();
printf( "Hello, World!\n" );
return 0;
}
編譯運行結果還是一樣的，可以看到CUDA C為標准C增加了__global__修飾符，這個修飾符告訴編譯器，函數應該編譯為在設備（Device）而不是主機(Host)上運行，CUDA把GPU稱為設備（Device），把CPU稱為主機（Host），而在GPU設備上執行的函數稱為核函數（Kernel），在設備上運行的函數需要增加__global__或__device__修飾符。
調用核函數增加了<<<1,1>>>修飾符，其它沒有任何變化。當然，這段代碼GPU並沒有做什么實際的工作，下面讓GPU實際做點事情。
GPU上的加法
修改上面的代碼：
__global__ void add( int a, int b, int *c )
{
*c = a + b;
}
int main( void )
{
int c;
int *dev_c;
cudaMalloc( &dev_c, sizeof(int) ) ;
add<<<1,1>>>( 2, 7, dev_c );
cudaMemcpy( &c, dev_c, sizeof(int),cudaMemcpyDeviceToHost );
printf( "2 + 7 = %d\n", c );
cudaFree( dev_c );
return 0;
}
主機調用核函數add做加法運算，函數add將在GPU上運行。需要注意的是，主機和設備在物理上處於不同的位置，使用了不同的內存，核函數不能直接使用主機上存儲的數據，同樣主機也不能直接使用設備上存儲的數據，數據需要在主機和設備之間傳輸。
在設備上分配內存使用cudaMalloc()，該函數類似malloc()。在執行完add函數后，計算結果保存在設備的內存上，還需要使用cudaMemcpy()傳輸到主機內存上，參數cudaMemcpyDeviceToHost表示從設備傳輸到主機。顯然，cudaMemcpyHostToDevice表示從主機傳輸到設備，cudaMemcpyDeviceToDevice表示從設備傳輸到另一個設備。
至此，一個完整地CUDA C代碼已經實現，下面來寫一段代碼查詢顯卡的設備參數。
查詢設備
在進行CUDA並行編程之前，對自己PC機的GPU設備性能及相關信息的了解是很有必要的，下面寫一段代碼來查詢設備參數信息。
查詢設備時會用的幾個函數：
cudaGetDeviceCount()，獲得CUDA設備的數量，一台PC可能會有多個CUDA設備，可以通過這個函數查詢。cudaGetDeviceProperties()，通過設備編號查詢設備屬性，設備編號從0開始。設備屬性保存在cudaDeviceProp結構體中，具體結構可查看cudaDeviceProp Struct Reference。
完整代碼如下：
int main(void)
{
cudaDeviceProp prop;
int count;
cudaGetDeviceCount(&count);
printf("cuda device count: %d\n", count);
for (int i = 0; i < count; ++i)
{
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 execiton timeout: " );
if (prop.kernelExecTimeoutEnabled )
{
printf ( "Enabled\n" );
}
else
{
printf ( "Disabled\n" );
}
printf ("integrated:");
if (prop.integrated)
{
printf("true\n");
}
else
{
printf("false\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 )", prop.maxGridSize[0], prop.maxGridSize[1], prop.maxGridSize[2] );
printf ("\n");
}
return 0;
}
運行這段代碼，就可以知道自己PC機配備GPU的具體信息，這對以后寫代碼是很重要的。
參考文獻
- 庫克. CUDA並行程序設計. 機械工業出版社, 2014.
- 桑德斯. GPU高性能編程CUDA實戰. 機械工業出版社, 2011.
- CUDA C Programming Guide
- CUDA Toolkit Documentation