更新( 2017 年 1 月):看看新的, 更容易介紹 CUDA !
本文是 CUDA C 和 C ++的一個系列,它是 CUDA 并行計算平臺的 C / C ++接口。本系列文章假定您熟悉 C 語言編程。我們將針對 Fortran 程序員運行一系列關于 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 及其存儲器。在主機上運行的代碼可以管理主機和設備上的內存,還可以啟動在設備上執行的函數 ?kernels 。這些內核由許多 GPU 線程并行執行。
鑒于 CUDA 編程模型的異構性, CUDA C 程序的典型操作序列是:
- 聲明并分配主機和設備內存。
- 初始化主機數據。
- 將數據從主機傳輸到設備。
- 執行一個或多個內核。
- 將結果從設備傳輸到主機。
記住這個操作序列,讓我們看一個 CUDA C 示例。
第一個 CUDA C 程序
在最近的一篇文章中,我演示了 薩克斯比的六種方法 ,其中包括一個 CUDA C 版本。 SAXPY 代表“單精度 A * X + Y ”,是并行計算的一個很好的“ hello world ”示例。在這篇文章中,我將剖析 CUDA C 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; 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 on 1M elements saxpy<<<(N+255)/256, 256>>>(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 ", maxError); cudaFree(d_x); cudaFree(d_y); free(x); free(y); }
函數 saxpy
是在 GPU 上并行運行的內核, main
函數是宿主代碼。讓我們從宿主代碼開始討論這個程序。
主機代碼
main 函數聲明兩對數組。
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
指向以典型方式使用malloc
分配的主機陣列,d_x
和d_y
數組指向從CUDA運行時API使用cudaMalloc
函數分配的設備數組。CUDA中的主機和設備有獨立的內存空間,這兩個空間都可以從主機代碼進行管理(CUDAC內核也可以在支持它的設備上分配設備內存)。
然后,主機代碼初始化主機數組。在這里,我們設置了一個 1 數組,以及一個 2 數組。
for (int i = 0; i < N; i++) { x[i] = 1.0f; y[i] = 2.0f; }
為了初始化設備數組,我們只需使用 cudaMemcpy
將數據從 x
和 y
復制到相應的設備數組 d_x
和 d_y
,它的工作方式與標準的 C memcpy
函數一樣,只是它采用了第四個參數,指定了復制的方向。在本例中,我們使用 cudaMemcpyHostToDevice
指定第一個(目標)參數是設備指針,第二個(源)參數是主機指針。
cudaMemcpy(d_x, x, N*sizeof(float), cudaMemcpyHostToDevice); cudaMemcpy(d_y, y, N*sizeof(float), cudaMemcpyHostToDevice);
在運行內核之后,為了將結果返回到主機,我們使用 cudaMemcpy
和 cudaMemcpyDeviceToHost
,從 d_y
指向的設備數組復制到 y
指向的主機數組。
cudaMemcpy(y, d_y, N*sizeof(float), cudaMemcpyDeviceToHost);
啟動內核
cord [EZX13 內核由以下語句啟動:
saxpy<<<(N+255)/256, 256>>>(N, 2.0, d_x, d_y);
三個 V 形符號之間的信息是 執行配置 ,它指示有多少設備線程并行執行內核。在 CUDA 中,軟件中有一個線程層次結構,它模仿線程處理器在 GPU 上的分組方式。在 CUDA 編程模型中,我們談到啟動一個 grid 為 螺紋塊 的內核。執行配置中的第一個參數指定網格中線程塊的數量,第二個參數指定線程塊中的線程數。
線程塊和網格可以通過為這些參數傳遞 dim3 (一個由 CUDA 用 x
、 y
和 z
成員定義的簡單結構)值來生成一維、二維或三維的線程塊和網格,但是對于這個簡單的示例,我們只需要一維,所以我們只傳遞整數。在本例中,我們使用包含 256 個線程的線程塊啟動內核,并使用整數算術來確定處理數組( (N+255)/256
)的所有 N
元素所需的線程塊數。
對于數組中的元素數不能被線程塊大小平均整除的情況,內核代碼必須檢查內存訪問是否越界。
清理
完成后,我們應該釋放所有分配的內存。對于使用 cudaMalloc()
分配的設備內存,只需調用 cudaFree()
。對于主機內存,請像往常一樣使用 free()
。
cudaFree(d_x); cudaFree(d_y); free(x); free(y);
設備代碼
現在我們繼續討論內核代碼。
__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__
de __global__
說明符定義諸如 Clara 這樣的內核。設備代碼中定義的變量不需要指定為設備變量,因為假定它們駐留在設備上。在這種情況下, n
、 a
和 i
變量將由每個線程存儲在寄存器中,指針 x
和 y
必須是指向設備內存地址空間的指針。這確實是真的,因為當我們從宿主代碼啟動內核時,我們將 d_x
和 d_y
傳遞給了內核。但是,前兩個參數 n
和 a
沒有在主機代碼中顯式傳輸到設備。因為函數參數在 C / C ++中是默認通過值傳遞的,所以 CUDA 運行時可以自動處理這些值到設備的傳輸。 CUDA 運行時 API 的這一特性使得在 GPU 上啟動內核變得非常自然和簡單——這幾乎與調用 C 函數一樣。
在我們的 saxpy
內核中只有兩行。如前所述,內核由多個線程并行執行。如果我們希望每個線程處理結果數組的一個元素,那么我們需要一種區分和標識每個線程的方法。 CUDA 定義變量 blockDim
、 blockIdx
和 threadIdx
。這些預定義變量的類型為 dim3
,類似于主機代碼中的執行配置參數。預定義變量 blockDim
包含在內核啟動的第二個執行配置參數中指定的每個線程塊的維度。預定義變量 threadIdx
和 blockIdx
分別包含線程塊中線程的索引和網格中的線程塊的索引。表達式:
int i = blockDim.x * blockIdx.x + threadIdx.x
生成用于訪問數組元素的全局索引。我們在這個例子中沒有使用它,但是還有一個 gridDim
,它包含在啟動的第一個執行配置參數中指定的網格維度。
在使用該索引訪問數組元素之前,將根據元素的數量 n
檢查其值,以確保沒有越界內存訪問。如果一個數組中的元素數不能被線程塊大小平均整除,并且結果內核啟動的線程數大于數組大小,則需要進行此檢查。內核的第二行執行 SAXPY 的元素級工作,除了邊界檢查之外,它與 SAXPY 主機實現的內部循環相同。
if (i < n) y[i] = a*x[i] + y[i];
編譯和運行代碼
CUDA C 編譯器 nvcc 是 NVIDIA CUDA 工具箱 的一部分。為了編譯我們的 SAXPY 示例,我們將代碼保存在一個擴展名為. cu 的文件中,比如說 saxpy.cu
。然后我們可以用 nvcc 編譯它。
nvcc -o saxpy saxpy.cu
然后我們可以運行代碼:
% ./saxpy Max error: 0.000000
總結與結論
通過對 SAXPY 的一個簡單的 CUDA C 實現的演練,您現在了解了編程 CUDA C 的基本知識。將 C 代碼“移植”到 CUDA C 只需要幾個 C 擴展:設備內核函數的 __global__
de Clara 說明符;啟動內核時使用的執行配置;內置的設備變量 blockDim
、 blockIdx
和 threadIdx
用來識別和區分并行執行內核的 GPU 線程。
異類 CUDA 編程模型的一個優點是,將現有代碼從 C 移植到 CUDA C 可以逐步完成,一次只能移植一個內核。
在本系列的下一篇文章中,我們將研究一些性能度量和度量。
注:本文基于 Gregory Reutsch 先生 的“ CUDA Fortran 簡介 ”一文。
?