前言

CUDA並行程式設計系列是本人在學習CUDA時整理的資料，內容大都來源於對《CUDA並行程式設計：GPU程式設計指南》、《GPU高效能程式設計CUDA實戰》和CUDA Toolkit Documentation的整理。通過本系列整體介紹CUDA並行程式設計。內容包括GPU簡介、CUDA簡介、環境搭建、執行緒模型、記憶體、原子操作、同步、流和多GPU架構等。

本系列目錄：

本文對CUDA進行簡單介紹，並通過例項程式碼演示怎麼編寫在GPU上執行的程式碼，最後寫一段程式碼來查詢本機GPU的裝置引數。

CUDA C簡介

CUDA C是NVIDIA公司設計的一種程式語言，用於在GPU上得編寫通用計算程式，目前也叫“ CUDA C and C++”，或者“CUDA C/C++”。CUDA C是C語言的擴充套件，所以使用起來和C語言類似。當然，CUDA現在已經不侷限於C語言了，在NVIDIA ZONE的

GPU計算能力

在檢視某個顯示卡是否支援CUDA時，還會看到一個計算能力（Compute Capability）的引數。正如不同的CPU有著不同的功能和指令集，對於支援CUDA的GPU也同樣如此。NVIDIA將GPU支援的各種功能統稱為計算能力。硬體的計算能力是固定的。不同計算能力具有一定差別，如，在計算能力1.0版本不支援全域性記憶體上的原子操作。更高計算能力的GPU是低計算能力的超級，所以計算能力2.0支援的功能在3.0也全部支援。可以看到，目前GPU最高的計算能力已經達到5.3（Tegra X1）。

環境搭建

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裝置效能及相關資訊的瞭解是很有必要的，下面寫一段程式碼來查詢裝置引數資訊。

查詢裝置時會用的幾個函式：

1. cudaGetDeviceCount()，獲得CUDA裝置的數量，一臺PC可能會有多個CUDA裝置，可以通過這個函式查詢。
2. 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的具體資訊，這對以後寫程式碼是很重要的。

參考文獻