程序師世界是廣大編程愛好者互助、分享、學習的平台,程序師世界有你更精彩!
首頁
編程語言
C語言|JAVA編程
Python編程
網頁編程
ASP編程|PHP編程
JSP編程
數據庫知識
MYSQL數據庫|SqlServer數據庫
Oracle數據庫|DB2數據庫
 程式師世界 >> 編程語言 >> C語言 >> 關於C語言 >> CUDA 2,CUDA

CUDA 2,CUDA

編輯:關於C語言

CUDA 2,CUDA


CUDA簡介

CUDA是並行計算的平台和類C編程模型,我們能很容易的實現並行算法,就像寫C代碼一樣。只要配備的NVIDIA GPU,就可以在許多設備上運行你的並行程序,無論是台式機、筆記本抑或平板電腦。熟悉C語言可以幫助你盡快掌握CUDA。

CUDA編程

CUDA編程允許你的程序執行在異構系統上,即CUP和GPU,二者有各自的存儲空間,並由PCI-Express 總線區分開。因此,我們應該先注意二者術語上的區分:

  • Host:CPU and itsmemory (host memory)
  • Device: GPU and its memory (device memory)

代碼中,一般用h_前綴表示host memory,d_表示device memory。

kernel是CUDA編程中的關鍵,他是跑在GPU的代碼,用標示符__global__注明。

host可以獨立於host進行大部分操作。當一個kernel啟動後,控制權會立刻返還給CPU來執行其他額外的任務。所以,CUDA編程是異步的。一個典型的CUDA程序包含由並行代碼補足的串行代碼,串行代碼由host執行,並行代碼在device中執行。host端代碼是標准C,device是CUDA C代碼。我們可以把所有代碼放到一個單獨的源文件,也可以使用多個文件或庫。NVIDIA C編譯器(nvcc)可以編譯host和device生成可執行程序。

這裡再次說明下CUDA程序的處理流程:

Memory操作

cuda程序將系統區分成host和device,二者有各自的memory。kernel可以操作device memory,為了能很好的控制device端內存,CUDA提供了幾個內存操作函數:

 

為了保證和易於學習,CUDA C 的風格跟C很接近,比如:

cudaError_t cudaMalloc ( void** devPtr, size_t size )

我們主要看看cudaMencpy,其函數原型為:

cudaError_t cudaMemcpy ( void* dst, const void* src, size_t count,cudaMemcpyKind kind )

其中cudaMemcpykind的可選類型有:

具體含義很好懂,就不多做解釋了。

對於返回類型cudaError_t,如果正確調用,則返回cudaSuccess,否則返回cudaErrorMemoryAllocation。可以使用char* cudaGetErrorString(cudaError_t error)將其轉化為易於理解的格式。

組織線程

掌握如何組織線程是CUDA編程的重要部分。CUDA線程分成Grid和Block兩個層次。

 

由一個單獨的kernel啟動的所有線程組成一個grid,grid中所有線程共享global memory。一個grid由許多block組成,block由許多線程組成,grid和block都可以是一維二維或者三維,上圖是一個二維grid和二維block。

這裡介紹幾個CUDA內置變量:

  • blockIdx:block的索引,blockIdx.x表示block的x坐標。
  • threadIdx:線程索引,同理blockIdx。
  • blockDim:block維度,上圖中blockDim.x=5.
  • gridDim:grid維度,同理blockDim。

一般會把grid組織成2D,block為3D。grid和block都使用dim3作為聲明,例如:

dim3 block(3);
// 後續博文會解釋為何這樣寫grid
dim3 grid((nElem+block.x-1)/block.x);

需要注意的是,dim3僅為host端可見,其對應的device端類型為uint3。

啟動CUDA kernel

CUDA kernel的調用格式為:

kernel_name<<<grid, block>>>(argument list);

其中grid和block即為上文中介紹的類型為dim3的變量。通過這兩個變量可以配置一個kernel的線程總和,以及線程的組織形式。例如:

kernel_name<<<4, 8>>>(argumentt list);

該行代碼表明有grid為一維,有4個block,block為一維,每個block有8個線程,故此共有4*8=32個線程。

注意,不同於c函數的調用,所有CUDA kernel的啟動都是異步的,當CUDA kernel被調用時,控制權會立即返回給CPU。

函數類型標示符

__device__ 和__host__可以組合使用。

kernel的限制:

  • 僅能獲取device memory 。
  • 必須返回void類型。
  • 不支持可變數目參數。
  • 不支持靜態變量。
  • 不支持函數指針。
  • 異步。

代碼分析

#include <cuda_runtime.h>
#include <stdio.h>
#define CHECK(call) \
{ \
  const cudaError_t error = call; \
  if (error != cudaSuccess) \
  { \
    printf("Error: %s:%d, ", __FILE__, __LINE__); \
    printf("code:%d, reason: %s\n", error, cudaGetErrorString(error)); \
    exit(1); \
  } \
}
void checkResult(float *hostRef, float *gpuRef, const int N) {   double epsilon = 1.0E-8;   bool match = 1;   for (int i=0; i<N; i++) {     if (abs(hostRef[i] - gpuRef[i]) > epsilon) {       match = 0;       printf("Arrays do not match!\n");       printf("host %5.2f gpu %5.2f at current %d\n",hostRef[i],gpuRef[i],i);       break;     }   }   if (match) printf("Arrays match.\n\n"); }
void initialData(float *ip,int size) {   // generate different seed for random number   time_t t;   srand((unsigned) time(&t));   for (int i=0; i<size; i++) {     ip[i] = (float)( rand() & 0xFF )/10.0f;   } }
void sumArraysOnHost(float *A, float *B, float *C, const int N) {   for (int idx=0; idx<N; idx++)   C[idx] = A[idx] + B[idx]; }
__global__ void sumArraysOnGPU(float *A, float *B, float *C) {   int i = threadIdx.x;   C[i] = A[i] + B[i]; }
int main(int argc, char **argv) {   printf("%s Starting...\n", argv[0]);   // set up device   int dev = 0;   cudaSetDevice(dev);
  // set up data size of vectors   int nElem = 32;   printf("Vector size %d\n", nElem);
  // malloc host memory   size_t nBytes = nElem * sizeof(float);   float *h_A, *h_B, *hostRef, *gpuRef;   h_A = (float *)malloc(nBytes);   h_B = (float *)malloc(nBytes);   hostRef = (float *)malloc(nBytes);   gpuRef = (float *)malloc(nBytes);
  // initialize data at host side   initialData(h_A, nElem);   initialData(h_B, nElem);   memset(hostRef, 0, nBytes);   memset(gpuRef, 0, nBytes);
  // malloc device global memory   float *d_A, *d_B, *d_C;   cudaMalloc((float**)&d_A, nBytes);   cudaMalloc((float**)&d_B, nBytes);   cudaMalloc((float**)&d_C, nBytes);
  // transfer data from host to device   cudaMemcpy(d_A, h_A, nBytes, cudaMemcpyHostToDevice);   cudaMemcpy(d_B, h_B, nBytes, cudaMemcpyHostToDevice);
  // invoke kernel at host side   dim3 block (nElem);   dim3 grid (nElem/block.x);   sumArraysOnGPU<<< grid, block >>>(d_A, d_B, d_C);   printf("Execution configuration <<<%d, %d>>>\n",grid.x,block.x);
  // copy kernel result back to host side   cudaMemcpy(gpuRef, d_C, nBytes, cudaMemcpyDeviceToHost);
  // add vector at host side for result checks   sumArraysOnHost(h_A, h_B, hostRef, nElem);
  // check device results   checkResult(hostRef, gpuRef, nElem);
  // free device global memory   cudaFree(d_A);   cudaFree(d_B);   cudaFree(d_C);
  // free host memory   free(h_A);   free(h_B);   free(hostRef);   free(gpuRef);   return(0); }

編譯指令:$nvcc sum.cu -o sum

運行: $./sum

輸出:

./sum Starting...
Vector size 32
Execution configuration <<<1, 32>>>
Arrays match.

代碼下載:CodeSamples.zip

  1. 上一頁:
  2. 下一頁:
Copyright © 程式師世界 All Rights Reserved