CUDA中文教程02之心得體會

第二集主要是講了CUDA程式設計模型，真正的介紹瞭如何編寫一個CUDA程式。

首先是介紹了一些基本的概念及資料型別。作為CPU的協處理器，GPU有自己的儲存器，可以並行的進行多執行緒計算，是一種並行的處理裝置。

接下來是介紹了一些關鍵字以及應用程式介面等。CUDA語言跟C語言相似。一個CUDA kernel 是用一個序列陣列來執行的，所有的執行緒都執行一樣的程式碼，每一個執行緒都有自己的ID，負責被指定的任務。

前面提過，一個block包含了512個threads，在同一個block內的thread可以通過shared memory 同步合作，但是不同的block裡的thread 是不能進行合作的。每個執行緒有自己ID去決定處理什麼資料。block也有自己的ID，現在的block ID只包含兩維，e.g. Block（0，0），Block （1，0）。thread ID包含三維，e.g. Thread （0，0，0）。

在CPU（host）和GPU（device）之間進行讀寫資料操作要通過 Global memory。它的內容對於所有執行緒來說是可見的。但是每一次訪問Global memory 有很長的延遲。

CUDA的應用程式藉口（API）是C語言的擴充套件。

CUDA的儲存分配函式cudaMalloc（），函式將物件分配到Global memory 中，這個函式需要兩個引數，1、指向這個物件的指標地址，2、分配物件的大小。當完成任務後，用cudaFree（）釋放Global memory裡面的物件。

從host到Global memory 的資料傳輸，用到的函式是cudaMemcpy（），這個函式需要四個引數，1、指向目的地的指標，2、指向源頭的指標，3、需要複製的位元組數，4、傳輸的型別（host to host / host to device / device to host / device to device）。

下面是一個例子：//從host到device的傳輸，cudaMemcpyHostToDevice是一個象徵常量

                        cudaMemcpy（Md.elements,M.elements,size,cudaMemcpyHostToDevice）;

                       //從device到host的傳輸，cudaMemcpyDeviceToHost是一個象徵常量

                        cudaMemcpy（M.elements,Md.elements,size,cudaMemcpyDeviceToHost）;

CUDA關鍵字。

1、當函式是在device上獲取且執行的，可以在函式前冠以__device__:

e.g.  __device__ float DeviceFunc()

2、當函式是在host上獲取但是在device上執行的，可以在函式前冠以__global__:
e.g. __global__ void KernelFunc() //must return void
3、當函式是在host上獲取且執行的，可以在函式前冠以__host__:
e.g.  __host__ float HostFunc()

一般情況下，__host__可以省略，是一種預設狀態,只有當定義在device上的函式也想要在host上使用，就可以同時將——__device__和__host__同時冠以函式前。

注意，在device上執行的函式沒有遞迴，沒有靜態變數，沒有可變引數。

矩陣相乘例子：

void MatrixMulOnDevice(float* M, float* N, float* P, int Width)
{
   int size = Width * Width * sizeof(float);
    float* Md, Nd, Pd;
   …
1. // Allocate and Load M, N to device memory
    cudaMalloc(&Md, size);
    cudaMemcpy(Md, M, size, cudaMemcpyHostToDevice);
     cudaMalloc(&Nd, size);
     cudaMemcpy(Nd, N, size, cudaMemcpyHostToDevice);
     // Allocate P on the device
    cudaMalloc(&Pd, size);
 
2.  // Kernel invocation code   
   
// Matrix multiplication kernel – per thread code
__global__ void MatrixMulKernel(float* Md, float* Nd, float* Pd, int Width)
{
    // 2D Thread ID
    int tx = threadIdx.x;
    int ty = threadIdx.y;
    // Pvalue is used to store the element of the matrix
    // that is computed by the thread
    float Pvalue = 0;   
    for (int k = 0; k < Width; ++k)
    {
         float Melement = Md[ty * Width + k];
         float Nelement = Nd[k * Width + tx];
         Pvalue += Melement * Nelement;
    }
    // Write the matrix to device memory;
    // each thread writes one element
    Pd[ty * Width + tx] = Pvalue;
}
 
    // Setup the execution configuration
    dim3 dimBlock(Width, Width);
    dim3 dimGrid(1, 1);
    // Launch the device computation threads!
    MatrixMulKernel<<>>(Md, Nd, Pd);
 
3.    // Read P from the device
      cudaMemcpy(P, Pd, size, cudaMemcpyDeviceToHost);
       // Free device matrices
      cudaFree(Md); cudaFree(Nd); cudaFree (Pd);
}

源自：xieyan blog