原文地址，純翻譯
https://developer.nvidia.com/blog/easy-introduction-cuda-c-and-c/
這是cuda並行計算平台 c和c++接口系列的第一篇文章。學習前要求熟練掌握c，針對cuda fortran編程的帖子也會同步更新。這兩個系列將涵蓋cuda平台上並行計算基本概念。從這里開始，除非有特別說明，否則我將使用屬於"cuda c"作為cuda c和c++的簡寫。cuda c本質上是帶有一些拓展的c/c++，允許使用多個並行線程在gpu上執行的函數。
CUDA編程模型基礎
在接觸cuda c代碼之前，那些剛接觸cuda的人最好先了解cuda編程模型基本描述和其中的一些術語。
cuda編程模型是一種同時使用cpu和gpu的異構模型。在cuda中，host指cpu及其內存，device指gpu及其內存，host上運行的代碼可以管理host和device上的內存、啟動kernel(核函數)，這些kernel是device上運行的函數，他們由gpu上的許多線程並發運行。
鑒於cuda編程模型的異構性質，cuda c程序的普遍操作順序是：
1.聲明和分配host端，device端內存
2.初始化host端數據
3.將數據從host端傳送到device端
4.執行一個或多個核函數
5.將結果從device端傳回host端
記住這個操作流程，讓我們來看一個cuda c的例子
第一個cuda c程序
上一篇文章中，我介紹了六種SAXPY(Scalar Alpha X Plus Y)的方法,其中就包括了cuda c版本，SAXPY表示單精度A*X+Y，對於並行計算來說是一個很好的hello world程序。在這篇文章中，我將展示一個cuda SAXPY的更完整版本，詳細說明做了什么以及為什么這樣做，完整的SAXPY代碼如下：
#include <stdio.h> __global__ void saxpy(int n,float a,float *x,float *y) { int i=blockIdx.x*blockDim.x+threadIdx.x; if(i<n) { y[i]=a*x[i]+y[i]; } } int main(void) { int N=1<<20; // 1左移20 float *x,*y,*d_x,*d_y; x=(float*)malloc(N*sizeof(float)); y=(float*)malloc(N*sizeof(float)); cudaMalloc(&d_x,N*sizeof(float)); cudaMalloc(&d_y,N*sizeof(float)); for(int i=0;i<N;++i) { x[i]=1.0f; y[i]=2.0f; } cudaMemcpy(d_x,x,N*sizeof(float),cudaMemcpyHostToDevice); cudaMemcpy(d_y,y,N*sizeof(float),cudaMemcpyHostToDevice); // perform SAXPY in 1M elements int threads=256; int blocks=(N+threads-1)/threads; // 上取整 saxpy<<<blocks,threads>>>(N,2.0f,d_x,d_y); cudaMemcpy(y,d_y,N*sizeof(float),cudaMemcpyDeviceToHost); float maxError=0.0f; for(int i=0;i<N;++i) { maxError=max(maxError,abs(y[i]-4.0f)); } printf("Max error: %f\n",maxError); cudaFree(d_x); cudaFree(d_y); free(x); free(y); return 0; }
函數saxpy就是在gpu上運行的核函數，main函數是host端代碼，讓我們先從host端代碼開始討論這個程序
Host端代碼
主函數聲明了兩對數組
float *x,*y,*d_x,*d_y; x=(float*)malloc(N*sizeof(float)); y=(float*)malloc(N*sizeof(float)); cudaMalloc(&d_x,N*sizeof(float)); cudaMalloc(&d_y,N*sizeof(float));
指針x和y指向host端數組，使用malloc函數申請內存空間，d_x,d_y數組指向device端數組，使用cuda runtime api中的cudaMalloc函數申請空間。cuda編程中host端和device端具有獨立的內存空間，兩者都可以通過host端代碼進行管理（cuda c內核還可以在支持他的設備上分配設備內存）。
host端代碼隨后初始化host端數組。在此，我們將x設為元素全部為1的數組，y設為元素全為2的數組
for(int i=0;i<N;++i) { x[i]=1.0f; y[i]=2.0f; }
要初始化device端數組，我們只需使用cudaMemcpy函數將數據從x，y拷貝到device端對應數組d_x,d_y上，這個過程就像c語言中的memcpy函數，唯一的區別就是cudaMemcpy需要地四個參數來指定數據拷貝到方向（host端到device端還是device端到host端），在此例中，我們使用cudaMemcpyHostToDevice表示數據從host拷貝到device端。
cudaMemcpy(d_x,x,N*sizeof(float),cudaMemcpyHostToDevice); cudaMemcpy(d_y,y,N*sizeof(float),cudaMemcpyHostToDevice);
在運行完核函數后，為了將數據d_y從device端拷貝回host端的y，我們使用cudaMemcpy函數，第四個參數指定為cudaMemcpyDeviceToHost
cudaMemcpy(y,d_y,N*sizeof(float),cudaMemcpyDeviceToHost);
啟動核函數
saxpy核函數由以下代碼啟動
saxpy<<<blocks,threads>>>(N,2.0f,d_x,d_y);
<<<>>>中間的信息是核函數的執行配置，具體指有多少設備線程並行執行內核。在cuda中，軟件有一個線程層次結構，它模仿線程處理器在gpu上的分組方式。在cuda編程模型中，我們通常說通過grid來啟動線程block。執行配置中的第一個參數指定grid中的block數目，第二個參數指定一個block中的線程數。
block和grid可以通過給dim3傳值構造成1維，2維或3維度數據結構（由cuda定義的具有x,y,z策划那個圓的簡單結構），但是對於這個簡單的例子，我們只需要一個維度，所以我們傳遞一個整數。在這個例子中，我們設置一個block中包含256個線程，使用取整算法來計算得到處理數組N個元素所需要的block數目（(N+256-1)/256）。
對於數組中元素的數量不能被線程塊大小整除的情況，內核代碼必須檢查越界內存訪問。（我認為就是檢測取整多出來的那部分）
釋放內存
程序結束時，我們需要釋放掉所有申請的內存空間。對於用cudaMalloc()申請的device端內存，使用cudaFree()來釋放。對於host端申請單內存，使用free()來釋放。
cudaFree(d_x); cudaFree(d_y); free(x); free(y);
device端代碼
現在我們來看核函數代碼
__global__ void saxpy(int n,float a,float *x,float *y) { int i=blockIdx.x*blockDim.x+threadIdx.x; if(i<n) { y[i]=a*x[i]+y[i]; } }
在cuda中，我們使用__global__聲明符來聲明核函數。device代碼中定義的變量不需要指定為device變量，因為他們被指定駐留在device上。（我的理解是核函數里聲明的變量，他的生命周期就是該線程上執行的核函數生命周期）。在本例中，n，a和i變量將由每個線程存儲在一個寄存器中，並且指針x和y必須是指向device內存地址空間的指針。因為當我們從host端啟動核函數時，我們要將d_x,d_y傳遞給kernel。前兩個變量n和a，我們沒有明確的代碼將它們從host端傳到device端，由於函數參數在c/c++中默認按值傳遞，因此cuda運行時可以自動處理這些值到device的傳輸。cuda runtime api這個特性使得gpu上啟動核函數非常自然和容易，幾乎與調用c函數一樣。
核函數saxpy的代碼只有兩行，前面我們提到過，核函數是由多線程並發執行的。如果我們希望每個線程處理數組的一個元素，那么我們需要一種區分和識別每個線程的方法。cuda定義了變量blockDim，blockIdx和threadIdx。這些預定義變量的類型為dim3，類似於主機代碼中的執行配置參數。預定義變量blockDim包含在核函數啟動的第二個執行配置參數中，指定block中的線程個數（維度）。預定義變量threadIdx和blockIdx分別表示block內的線程索引和grid內地block索引。表達方式：
int i=blockIdx.x*blockDim.x+threadIdx.x;
生成用於訪問數組元素的全局索引。我們在這個例子中沒有使用gridDim，gridDim表示核函數啟動的第一個執行配置參數中指定的grid尺寸（一個grid里面有多少個block）
在使用索引訪問數組元素之前，會根據元素個數n檢查索引是否合法，以確保沒有內存越界訪問。如果數組中的元素不能被block中的線程數整除，（由於向上取整）核函數啟動的線程數會大於數組大小，此時需要進行檢查。核函數的第二行逐個元素執行SAXPY操作，除了邊界檢查外，他的結果等於host端通過循環實現的SAXPY
if(i<n) y[i]=a*x[i]+y[i];
編譯和運行代碼
cuda c編譯器，nvcc是NVIDIA CUDA Toolkit（http://www.nvidia.com/content/cuda/cuda-toolkit.html）的一部分。為了編譯我們的SAXPY例子，我們將代碼保存為.cu格式，文件名是saxpy.cu。我們可以使用nvcc編譯它
nvcc -o saxpy saxpy.cu
我們可以運行代碼
./saxpy
總結與結論
通過使用cuda c對saxpy簡單實現，我們現在了解了cuda編程的基礎知識。將c代碼移植到cuda c上只需要幾個c擴展：核函數的__global__聲明符；核函數啟動時的執行配置（<<<blocks,threads>>>）；用於識別和區分並行執行核函數的gpu線程的內置設備變量blockDim,blockIdx和threadIdx。
異構cuda編程模型的一個優點是，可以增量的將現有代碼從c移植到cuda c，一次一個內核。
在本系列的下一篇文章中，我么將了解一些性能測量和指標。