• <xmp id="om0om">
  • <table id="om0om"><noscript id="om0om"></noscript></table>
  • 高性能計算

    CUDA C 和 C++ 的簡單介紹

    ?
    ?

    CUDA Cube更新( 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 程序的典型操作序列是:

    1. 聲明并分配主機和設備內存。
    2. 初始化主機數據。
    3. 將數據從主機傳輸到設備。
    4. 執行一個或多個內核。
    5. 將結果從設備傳輸到主機。

    記住這個操作序列,讓我們看一個 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));

    指針xy指向以典型方式使用malloc分配的主機陣列,d_xd_y數組指向從CUDA運行時API使用cudaMalloc函數分配的設備數組。CUDA中的主機和設備有獨立的內存空間,這兩個空間都可以從主機代碼進行管理(CUDAC內核也可以在支持它的設備上分配設備內存)。

    然后,主機代碼初始化主機數組。在這里,我們設置了一個 1 數組,以及一個 2 數組。

      for (int i = 0; i < N; i++) {
        x[i] = 1.0f;
        y[i] = 2.0f;
      }

    為了初始化設備數組,我們只需使用 cudaMemcpy 將數據從 xy 復制到相應的設備數組 d_xd_y ,它的工作方式與標準的 C memcpy 函數一樣,只是它采用了第四個參數,指定了復制的方向。在本例中,我們使用 cudaMemcpyHostToDevice 指定第一個(目標)參數是設備指針,第二個(源)參數是主機指針。

      cudaMemcpy(d_x, x, N*sizeof(float), cudaMemcpyHostToDevice);
      cudaMemcpy(d_y, y, N*sizeof(float), cudaMemcpyHostToDevice);

    在運行內核之后,為了將結果返回到主機,我們使用 cudaMemcpycudaMemcpyDeviceToHost ,從 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 用 xyz 成員定義的簡單結構)值來生成一維、二維或三維的線程塊和網格,但是對于這個簡單的示例,我們只需要一維,所以我們只傳遞整數。在本例中,我們使用包含 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 這樣的內核。設備代碼中定義的變量不需要指定為設備變量,因為假定它們駐留在設備上。在這種情況下, nai 變量將由每個線程存儲在寄存器中,指針 xy 必須是指向設備內存地址空間的指針。這確實是真的,因為當我們從宿主代碼啟動內核時,我們將 d_xd_y 傳遞給了內核。但是,前兩個參數 na 沒有在主機代碼中顯式傳輸到設備。因為函數參數在 C / C ++中是默認通過值傳遞的,所以 CUDA 運行時可以自動處理這些值到設備的傳輸。 CUDA 運行時 API 的這一特性使得在 GPU 上啟動內核變得非常自然和簡單——這幾乎與調用 C 函數一樣。

    在我們的 saxpy 內核中只有兩行。如前所述,內核由多個線程并行執行。如果我們希望每個線程處理結果數組的一個元素,那么我們需要一種區分和標識每個線程的方法。 CUDA 定義變量 blockDimblockIdxthreadIdx 。這些預定義變量的類型為 dim3 ,類似于主機代碼中的執行配置參數。預定義變量 blockDim 包含在內核啟動的第二個執行配置參數中指定的每個線程塊的維度。預定義變量 threadIdxblockIdx 分別包含線程塊中線程的索引和網格中的線程塊的索引。表達式:

        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 說明符;啟動內核時使用的執行配置;內置的設備變量 blockDimblockIdxthreadIdx 用來識別和區分并行執行內核的 GPU 線程。

    異類 CUDA 編程模型的一個優點是,將現有代碼從 C 移植到 CUDA C 可以逐步完成,一次只能移植一個內核。

    在本系列的下一篇文章中,我們將研究一些性能度量和度量。

    注:本文基于 Gregory Reutsch 先生 的“ CUDA Fortran 簡介 ”一文。

    ?

    ?

    0

    標簽

    人人超碰97caoporen国产