CUDA入門教程

1. 准備makefile

為了避免每次都要鍵入nvcc的命令，要准備一個makefile。makefile如下：

CUFLAG = -g  -Xcompiler -v \
        -gencode=arch=compute_20,code=sm_20\
        -gencode=arch=compute_20,code=compute_20\
        -O2
IFLAGS = -I$(CUDA_DIR)/include -I$(CUDA_SDK_DIR)/C/common/inc -I../include
LFLAGS = -L$(CUDA_DIR)/lib64 -L$(CUDA_SDK_DIR)/C/lib
PRG = cuda_test
$(PRG) : main.cu
 nvcc main.cu -o $(PRG) $(CUFLAG) $(IFLAGS) $(LFLAGS)

Ubuntu 12.04 下 CUDA 編程 http://www.linuxidc.com/linux/2014-06/103056.htm

Ubuntu 12.04 安裝 CUDA-5.5  http://www.linuxidc.com/Linux/2013-10/91101.htm

Ubuntu 11.10 上安裝CUDA開發環境 http://www.linuxidc.com/Linux/2012-04/58913.htm

Fedora 15系統下配置CUDA環境 http://www.linuxidc.com/Linux/2011-12/49874.htm

Ubuntu 11.04 安裝 NVIDIA CUDA 4.0 RC2 http://www.linuxidc.com/Linux/2011-10/46304.htm

Linux Mint 13/Ubuntu 12.04 配置CUDA 4.2 & OpenCV 2.4.2 方法 http://www.linuxidc.com/Linux/2013-10/91102.htm

2 異構計算(Heterogeneous Computing)

以下為幾個技術名詞的簡單介紹：

 

• 主機(host)：CPU及其內存(host memory)。
• 設備(device)：GPU及其內存(device memory)。
• 主機代碼(host code)：運行在CPU上的（一般來說「串行執行」的）代碼。
• 設備代碼(device code)：運行在GPU上的並行執行的代碼。
• 異構計算：由主機代碼(host code)和設備代碼(device code)協同執行完成的計算。

宏觀上看，GPU執行代碼的流程如下：

 

 

1. 將輸入數據通過PCI總線從CPU內存拷貝到GPU的DRAM中。
2. 從內存中加載需要執行的代碼到GPU後。
3. 數據和指令都就緒後，就可以執行了。注意，在執行的過程中，GPU會在片上緩存數據以提升性能。
4. 計算完畢後，將結果從GPU的DRAM中拷回CPU的Memory中。

例1: Hello World

#include<stdio.h>
#include<stdlib.h>
#include<cuda.h>
#include<cutil.h>

__global__ void mykernel(void) {
}

int main(void) {
 mykernel<<<1,1>>>();
 printf("Hello World!\n");
 return 0;
}

上述代碼編譯後運行生成可執行文件cuda_test，運行cuda_test後將輸出：

 

Hello World!

注意：

1. 調用kernel時需要三個尖括號
2. 包含必要的頭文件

CUDA C/C++中引入的新關鍵字__global__所修飾的函數有以下兩方面含義：

• 此函數代碼由設備執行
• 此函數由主機代碼調用

nvcc將源代碼分為設備函數和主機函數兩大類：

• 設備函數由NVIDA編譯器編譯
• 主機函數由主機上配置的編譯器編譯

三個尖括號標志著一個從主機代碼調用設備代碼的函數，稱為“啟動內核”(kernel launch)

 

例2: 整數相加

#include<stdio.h>
#include<stdlib.h>
#include<cuda.h>
#include<cutil.h>

__global__ void integer_add(int * a, int * b, int * c) {
 *c = *a + *b;
}

int main(void) {
 int a,b,c;
 int * d_a, * d_b, * d_c;
 int size = sizeof(int);
 cudaMalloc((void**)&d_a,size);
 cudaMalloc((void**)&d_b,size);
 cudaMalloc((void**)&d_c,size);
 printf("Enter two integers with a space to separate them:\n");
 scanf("%d %d",&a,&b);
 cudaMemcpy(d_a,&a,size,cudaMemcpyHostToDevice);
 cudaMemcpy(d_b,&b,size,cudaMemcpyHostToDevice);
 integer_add<<<1,1>>>(d_a,d_b,d_c);
 cudaMemcpy(&c,d_c,size,cudaMemcpyDeviceToHost);
 cudaFree(d_a);
 cudaFree(d_b);
 cudaFree(d_c);
 printf("Sum is %d\n",c);
 return 0;
}

__global__修飾的integer_add函數說明：

• integer_add函數將在設備上執行
• integer_add函數將被主機調用

由於integer_add函數在設備上執行，所以指針a，b，c應該指向設備內存。這說明需要在設備內存中為變量開辟內存。   設備內存和主機內存在物理上是完全分開的不同電子部件：

• 設備指針指向GPU內存的某個位置。設備指針可以從主機端傳給設備端或者從設備端傳給主機端，但是設備指針不能在主機端解引用。
• 主機指針指向CPU內存的某個位置。主機指針可以從設備端傳給主機端或者從主機端傳給設備端，但是主機指針不能在設備端解引用。

CUDA API提供的用於處理設備內存的函數有cudaMalloc, cudaFree, cudaMemcpy。語義上分別對應於C語言的malloc, free, memcpy函數。這幾個函數的具體使用方法如例2所示。  

3 塊(Blocks)

GPU是用來實現大規模並行的，如何實現呢？將上述例子擴展一下，如果我們要實現兩個向量相加： add<<<1,1>>>() ---> add<<<N,1>>> N表示同時調用N次add函數，這樣就可以實現並行的向量相加了。 每個被並行調用的add函數稱之為一個塊(block)。

• 塊的集合稱之為網格(grid).
• 每個塊可以使用索引值blockIdx.x

通過使用blockIdx.x作為索引，每個塊可以處理數組元素中的一部分。 有了這些基礎後，就可以實現並行版本的向量相加了。 例3:向量相加

#include<stdio.h>
#include<stdlib.h>
#include<cuda.h>
#include<cutil.h>
#include<time.h>

#define N 512

__global__ void vec_block_add(int * a, int * b, int * c) {
 c[blockIdx.x] = a[blockIdx.x] + b[blockIdx.x];
}

void rand_ints(int * arr, int count) {
 srand(time(NULL));
 for(int i=0;i<count;i++) {
  arr[i] = rand() % 100;
 }
}

int main(void) {
 int * a,* b,* c;
 int * d_a, * d_b, * d_c;
 int size = N * sizeof(int);
 cudaMalloc((void**)&d_a,size);
 cudaMalloc((void**)&d_b,size);
 cudaMalloc((void**)&d_c,size);
 
 a = (int *) malloc(size);
 rand_ints(a,N);
 b = (int *) malloc(size);
 rand_ints(b,N);
 c = (int *) malloc(size);
 
 cudaMemcpy(d_a,a,size,cudaMemcpyHostToDevice);
 cudaMemcpy(d_b,b,size,cudaMemcpyHostToDevice);
 vec_block_add<<<N,1>>>(d_a,d_b,d_c);
 cudaMemcpy(c,d_c,size,cudaMemcpyDeviceToHost);
 
#if 1
 for(int i=0;i<N;i++) {
  printf("%-5d: a:%-5d b:%-5d c:%-5d\n",i,a[i],b[i],c[i]);
 }
#endif
 
 cudaFree(d_a);
 cudaFree(d_b);
 cudaFree(d_c);
 
 free(a);
 free(b);
 free(c);
 return 0;
}

例3中最關鍵的代碼為如下幾行：

__global__ void vec_block_add(int * a, int * b, int * c) {
 c[blockIdx.x] = a[blockIdx.x] + b[blockIdx.x];
}

由於函數是並行執行的，和傳統的串行程序在integer_add函數中使用循環來完成加法相比，相當於由GPU這個加速器使用硬件的方式進行了循環展開，展開後便可以並行執行了。所以在編寫這段代碼時，需要使用blockIdx.x來定位當前執行的是循環的哪個部分。

從硬件的角度看，相當於同時有多個塊在並行執行： 塊0: c[0]=a[0]+b[0] 塊1: c[1]=a[1]+b[1] 塊2: c[2]=a[2]+b[2] 塊3: c[3]=a[3]+b[3] ....

更多詳情見請繼續閱讀下一頁的精彩內容： http://www.linuxidc.com/Linux/2014-07/104328p2.htm