2023年的深度學(xué)習(xí)入門指南(9) - SIMD和通用GPU編程

深度學(xué)習(xí)從一開(kāi)始就跟GPU有不解之緣，因?yàn)樗懔κ巧疃葘W(xué)習(xí)不可或缺的一部分。

時(shí)至今日，雖然多任務(wù)編程早已經(jīng)深入人心，但是很多同學(xué)還沒(méi)有接觸過(guò)CPU上的SIMD指令，更不用說(shuō)GPGPU的編程。這一篇我們先給SIMD和GPU編程掃個(gè)盲，讓大家以后用到的時(shí)候有個(gè)感性認(rèn)識(shí)。

CPU世界

從多線程說(shuō)起

曾經(jīng)的編程語(yǔ)言是不支持多線程的，需要操作系統(tǒng)和庫(kù)來(lái)提供多線程能力，比如pthread庫(kù)。時(shí)至今日，默認(rèn)不支持多線程的平臺(tái)還是有的，比如wasm。
1995年問(wèn)世的Java語(yǔ)言從1.0開(kāi)始就支持多線程，雖然一直到5.0版本才對(duì)多線程有重大改進(jìn)。C++語(yǔ)言從C++11開(kāi)始語(yǔ)言支持多線程了。

我們來(lái)看一個(gè)用C++多線程來(lái)實(shí)現(xiàn)矩陣乘法的例子：

#include <mutex>
    #include <thread>

        // 矩陣維度
        const int width = 4;

        // 矩陣
        int A[width][width] = {
            {1, 2, 3, 4},
            {5, 6, 7, 8},
            {9, 10, 11, 12},
            {13, 14, 15, 16}
        };
        int B[width][width] = {
            {1, 0, 0, 0},
            {0, 1, 0, 0},
            {0, 0, 1, 0},
            {0, 0, 0, 1}
        };
        int C[width][width] = {0};

        // 互斥鎖
        std::mutex mtx;

        // 計(jì)算線程
        void calculate(int row) {
            for (int col = 0; col < width; col++) {
                if (row < width && col < width) {
                    mtx.lock();
                    C[row][col] = A[row][col] + B[row][col];
                    mtx.unlock();
                }
            }
        }

        int main() {
            // 創(chuàng)建線程
            std::thread t1(calculate, 0);
            std::thread t2(calculate, 1);
            std::thread t3(calculate, 2);
            std::thread t4(calculate, 3);

            // 等待線程結(jié)束
            t1.join();
            t2.join();
            t3.join();
            t4.join();

            // 打印結(jié)果
            for (int i = 0; i < width; i++) {
                for (int j = 0; j < width; j++) {
                    printf("%d ", C[i][j]);
                }
                printf("\n");
            }
        }

我們給它配上一個(gè)CMakeLists.txt:

cmake_minimum_required(VERSION 3.10)

# Set the project name
project(MatrixAddO)

# Set the C++ standard
set(CMAKE_CXX_STANDARD 11)
set(CMAKE_CXX_STANDARD_REQUIRED True)

# Add the executable
add_executable(matrix_add matadd.cpp)

這個(gè)代碼大家應(yīng)該都比較熟悉，就不多解釋了?，F(xiàn)在支持C++11以上已經(jīng)是標(biāo)配了。

OpenMP

早在線程寫進(jìn)C++11標(biāo)準(zhǔn)之前，就有很多并發(fā)編程的框架了，比如MPI和OpenMP.

OpenMP是一套支持跨平臺(tái)共享內(nèi)存方式的多線程并發(fā)的編程API，使用C, C++和Fortran語(yǔ)言，可以在多種處理器體系和操作系統(tǒng)中運(yùn)行。它由OpenMP Architecture Review Board (ARB)牽頭提出，并由多家計(jì)算機(jī)硬件和軟件廠商共同定義和管理。

OpenMP最早是1997年發(fā)布的，當(dāng)時(shí)只支持Fortran語(yǔ)言。1998年開(kāi)始支持C/C++.

我們來(lái)看看用OpenMP如何實(shí)現(xiàn)矩陣的并發(fā)計(jì)算：

#include <iostream>
#include <omp.h>
#include <vector>

std::vector<std::vector<int>>
matrixAdd(const std::vector<std::vector<int>> &A,
          const std::vector<std::vector<int>> &B) {
  int rows = A.size();
  int cols = A[0].size();

  std::vector<std::vector<int>> C(rows, std::vector<int>(cols));

#pragma omp parallel for collapse(2)
  for (int i = 0; i < rows; i++) {
    for (int j = 0; j < cols; j++) {
      C[i][j] = A[i][j] + B[i][j];
    }
  }

  return C;
}

int main() {
  std::vector<std::vector<int>> A = {{1, 2, 3}, {4, 5, 6}, {7, 8, 9}};

  std::vector<std::vector<int>> B = {{9, 8, 7}, {6, 5, 4}, {3, 2, 1}};

  std::vector<std::vector<int>> C = matrixAdd(A, B);

  for (const auto &row : C) {
    for (int val : row) {
      std::cout << val << " ";
    }
    std::cout << std::endl;
  }

  return 0;
}                                 

#pragma omp parallel for collapse(2) 是一個(gè) OpenMP 編譯指令，用于表示一個(gè)并行區(qū)域，其中嵌套的循環(huán)將并行執(zhí)行。讓我們?cè)敿?xì)解釋這個(gè)指令的各個(gè)部分：

#pragma omp：這是一個(gè)編譯指令，表示接下來(lái)的代碼將使用 OpenMP 進(jìn)行并行化。

parallel for：這是一個(gè)組合指令，表示接下來(lái)的 for 循環(huán)將在多個(gè)線程上并行執(zhí)行。每個(gè)線程將處理循環(huán)的一部分，從而加速整個(gè)循環(huán)的執(zhí)行。

collapse(2)：這是一個(gè)可選子句，用于指示嵌套循環(huán)的并行化。在這個(gè)例子中，collapse(2) 表示將兩層嵌套的循環(huán)（即外層和內(nèi)層循環(huán)）合并為一個(gè)并行循環(huán)。這樣可以更好地利用多核處理器的性能，因?yàn)椴⑿卸仍黾恿恕?/p>

在我們的矩陣加法示例中，#pragma omp parallel for collapse(2) 指令應(yīng)用于兩個(gè)嵌套的 for 循環(huán)，它們分別遍歷矩陣的行和列。使用此指令，這兩個(gè)循環(huán)將合并為一個(gè)并行循環(huán)，從而在多核處理器上實(shí)現(xiàn)更高的性能。

需要注意的是，為了在程序中使用 OpenMP，你需要使用支持 OpenMP 的編譯器（如 GCC 或 Clang），并在編譯時(shí)啟用 OpenMP 支持（如在 GCC 中使用 -fopenmp 標(biāo)志）。

我們來(lái)寫個(gè)支持OpenMP的CMakeLists.txt:

cmake_minimum_required(VERSION 3.10)

# Set the project name
project(MatrixAddOpenMP)

# Set the C++ standard
set(CMAKE_CXX_STANDARD 11)
set(CMAKE_CXX_STANDARD_REQUIRED True)

# Find OpenMP
find_package(OpenMP REQUIRED)

# Add the executable
add_executable(matrix_add main.cpp)

# Link OpenMP to the executable
if(OpenMP_CXX_FOUND)
    target_link_libraries(matrix_add PUBLIC OpenMP::OpenMP_CXX)
endif()

可見(jiàn)，用了OpenMP的for循環(huán)，就可以變串行為并行。從而大大簡(jiǎn)化并行編程的難度。

SIMD

雖然多線程和OpenMP看起來(lái)都不錯(cuò)，都容易編程，但是，我們的優(yōu)化并不是以簡(jiǎn)化編程為目的的。

雖然我們抱怨Intel是牙膏廠，每年的進(jìn)步越來(lái)越有限。不過(guò)，還總是有新的指令增加到新的架構(gòu)中來(lái)。這其中就有越來(lái)越強(qiáng)大的SIMD指令。

SIMD就是一條機(jī)器指令可以實(shí)現(xiàn)多條數(shù)據(jù)的操作。在Intel平臺(tái)上，早在1997年就推出了64位的MMX指令集。1999年又有了128位的SSE指令集。2011年，又推出了256位的AVX(Advanced Vector Extensions)指令，我們來(lái)個(gè)例子看看：

#include <iostream>
#include <immintrin.h> // 包含 AVX 指令集頭文件

void matrix_addition_avx(float* A, float* B, float* C, int size) {
    for (int i = 0; i < size; i++) {
        for (int j = 0; j < size; j += 8) { // 每次處理 8 個(gè)元素（AVX 可以處理 256 位，即 8 個(gè)單精度浮點(diǎn)數(shù)）
            __m256 vecA = _mm256_loadu_ps(&A[i * size + j]);
            __m256 vecB = _mm256_loadu_ps(&B[i * size + j]);
            __m256 vecC = _mm256_add_ps(vecA, vecB);
            _mm256_storeu_ps(&C[i * size + j], vecC);
        }
    }
}

int main() {
    int size = 8; // 假設(shè)矩陣大小為 8x8
    float A[64] = { /* ... */ }; // 初始化矩陣 A
    float B[64] = { /* ... */ }; // 初始化矩陣 B
    float C[64] = { 0 }; // 結(jié)果矩陣 C

    matrix_addition_avx(A, B, C, size);

    // 輸出結(jié)果
    for (int i = 0; i < size; i++) {
        for (int j = 0; j < size; j++) {
            std::cout << C[i * size + j] << " ";
        }
        std::cout << std::endl;
    }

    return 0;
}

我們來(lái)解釋一下使用SIMD的幾條語(yǔ)句：

__m256 vecA = _mm256_loadu_ps(&A[i * size + j])：從矩陣 A 中加載 8 個(gè)浮點(diǎn)數(shù)（一次性處理 256 位數(shù)據(jù)），存儲(chǔ)在一個(gè)名為 vecA 的 __m256 類型變量中。

__m256 vecB = _mm256_loadu_ps(&B[i * size + j])：同樣地，從矩陣 B 中加載 8 個(gè)浮點(diǎn)數(shù)，存儲(chǔ)在一個(gè)名為 vecB 的 __m256 類型變量中。

__m256 vecC = _mm256_add_ps(vecA, vecB)：使用 AVX 指令 _mm256_add_ps 對(duì) vecA 和 vecB 中的浮點(diǎn)數(shù)分別進(jìn)行逐元素加法，并將結(jié)果存儲(chǔ)在名為 vecC 的 __m256 類型變量中。

_mm256_storeu_ps(&C[i * size + j], vecC)：將 vecC 中的 8 個(gè)加法結(jié)果存儲(chǔ)回矩陣 C 的相應(yīng)位置。

這段代碼使用了 AVX 指令集，實(shí)現(xiàn)了對(duì)浮點(diǎn)矩陣的加法運(yùn)算。請(qǐng)注意，為了充分利用 AVX 的并行處理能力，矩陣尺寸應(yīng)該是 8 的倍數(shù)。如果矩陣尺寸不是 8 的倍數(shù)，需要添加額外的邏輯來(lái)處理剩余的元素。

后來(lái)，Intel又推出了AVX2指令集，不過(guò)對(duì)于我們上邊的代碼并沒(méi)有太多優(yōu)化，而主要優(yōu)化是在整數(shù)方面。

上節(jié)我們學(xué)習(xí)的量化和解量化就用上了，我們這次使用AVX2提供的整數(shù)計(jì)算的加速來(lái)實(shí)現(xiàn)：

#include <iostream>
#include <immintrin.h> // 包含 AVX2 指令集頭文件

void matrix_addition_avx2_int(int *A, int *B, int *C, int size) {
    for (int i = 0; i < size; i++) {
        for (int j = 0; j < size; j += 8) { // 每次處理 8 個(gè)元素（AVX2 可以處理 256 位，即 8 個(gè) int32 整數(shù)）
            __m256i vecA = _mm256_loadu_si256((__m256i *)&A[i * size + j]);
            __m256i vecB = _mm256_loadu_si256((__m256i *)&B[i * size + j]);
            __m256i vecC = _mm256_add_epi32(vecA, vecB);
            _mm256_storeu_si256((__m256i *)&C[i * size + j], vecC);
        }
    }
}

int main() {
    int size = 8; // 假設(shè)矩陣大小為 8x8
    int A[64] = { /* ... */ }; // 初始化矩陣 A
    int B[64] = { /* ... */ }; // 初始化矩陣 B
    int C[64] = {0}; // 結(jié)果矩陣 C

    matrix_addition_avx2_int(A, B, C, size);

    // 輸出結(jié)果
    for (int i = 0; i < size; i++) {
        for (int j = 0; j < size; j++) {
            std::cout << C[i * size + j] << " ";
        }
        std::cout << std::endl;
    }

    return 0;
}

我們不惜折騰量化一把轉(zhuǎn)換成整數(shù)的原因是，AVX中只有_mm_add_epi32指令，只能對(duì)兩個(gè)128位整數(shù)向量的逐元素相加，而_mm256_add_epi32是256位，數(shù)據(jù)量加倍了。
不只是加法，AVX2 提供了一系列針對(duì)整數(shù)操作的新指令，例如乘法、位操作和打包/解包操作等。
AVX2指令的執(zhí)行吞吐量(throughput)一般為1指令/周期,而AVX1為2指令/周期。所以在同頻率下,AVX2的整數(shù)加法指令性能理論上可以提高一倍。
同時(shí)， 與其他AVX2指令結(jié)合使用，如_mm256_load_si256或_mm256_store_si256等，來(lái)從內(nèi)存中加載或存儲(chǔ)向量，這樣可以提高內(nèi)存訪問(wèn)的性能和帶寬。

后來(lái)，Intel還推出了AVX512指令，基本上就把AVX1中的256換成512就可以了：

#include <iostream>
#include <immintrin.h> // 包含 AVX-512 指令集頭文件

void matrix_addition_avx512(float *A, float *B, float *C, int size) {
    for (int i = 0; i < size; i++) {
        for (int j = 0; j < size; j += 16) { // 每次處理 16 個(gè)元素（AVX-512 可以處理 512 位，即 16 個(gè)單精度浮點(diǎn)數(shù)）
            __m512 vecA = _mm512_loadu_ps(&A[i * size + j]);
            __m512 vecB = _mm512_loadu_ps(&B[i * size + j]);
            __m512 vecC = _mm512_add_ps(vecA, vecB);
            _mm512_storeu_ps(&C[i * size + j], vecC);
        }
    }
}

int main() {
    int size = 16; // 假設(shè)矩陣大小為 16x16
    float A[256] = { /* ... */ }; // 初始化矩陣 A
    float B[256] = { /* ... */ }; // 初始化矩陣 B
    float C[256] = {0}; // 結(jié)果矩陣 C

    matrix_addition_avx512(A, B, C, size);

    // 輸出結(jié)果
    for (int i = 0; i < size; i++) {
        for (int j = 0; j < size; j++) {
            std::cout << C[i * size + j] << " ";
        }
        std::cout << std::endl;
    }

    return 0;
}

但是，優(yōu)化并不總是一根筋地往上堆指令就可以的，AVX512是一種非常耗電的指令集，此時(shí)我們需要實(shí)測(cè)權(quán)衡一下。

針對(duì)手機(jī)上用的ARM CPU，可以使用NEON指令來(lái)實(shí)現(xiàn)SIMD功能：

#include <stdio.h>
#include <arm_neon.h>

void matrix_addition_neon(float *A, float *B, float *C, int size) {
    for (int i = 0; i < size; i++) {
        for (int j = 0; j < size; j += 4) { // 每次處理 4 個(gè)元素（NEON 可以處理 128 位，即 4 個(gè)單精度浮點(diǎn)數(shù)）
            float32x4_t vecA = vld1q_f32(&A[i * size + j]);
            float32x4_t vecB = vld1q_f32(&B[i * size + j]);
            float32x4_t vecC = vaddq_f32(vecA, vecB);
            vst1q_f32(&C[i * size + j], vecC);
        }
    }
}

int main() {
    int size = 4; // 假設(shè)矩陣大小為 4x4
    float A[16] = { /* ... */ }; // 初始化矩陣 A
    float B[16] = { /* ... */ }; // 初始化矩陣 B
    float C[16] = {0}; // 結(jié)果矩陣 C

    matrix_addition_neon(A, B, C, size);

    // 輸出結(jié)果
    for (int i = 0; i < size; i++) {
        for (int j = 0; j < size; j++) {
            printf("%f ", C[i * size + j]);
        }
        printf("\n");
    }

    return 0;
}

對(duì)于初接觸匯編級(jí)優(yōu)化的同學(xué)，可能感覺(jué)很新鮮。不過(guò)，挑戰(zhàn)更大的在后面，我們要進(jìn)入GPU的世界了。

GPU世界

歡迎來(lái)到異構(gòu)計(jì)算的世界。之前我們的代碼不管怎么寫，都是在CPU上運(yùn)行的。
從這一時(shí)刻開(kāi)始，不管什么技術(shù)，我們都是由CPU和GPU兩部分代碼共同組合的了。

我們先從目前看仍然是主力的CUDA開(kāi)始。

CUDA

CUDA 1.0于2007年發(fā)布。目前CUDA版本為12.1。

目前廣泛適配的是CUDA 11.x，現(xiàn)在較新的版本為CUDA 11.8。因?yàn)镃UDA 11.x才支持A100為代表的安培架構(gòu)的GPU。3060，3070，3080，3090也是安培架構(gòu)的GPU。

2080, 2060, 1660這一系列的是圖靈架構(gòu)，對(duì)應(yīng)的是CUDA 10.x版本。

1060，1080這一系列對(duì)應(yīng)的是帕斯卡架我，對(duì)應(yīng)的是CUDA 8.0版本。

在CUDA中，運(yùn)行在GPU上的代碼我們叫做核函數(shù)。
我們先完整地看下這個(gè)代碼，然后再解釋。

#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <iostream>

// 矩陣加法的CUDA核函數(shù)
__global__ void matrixAdd10(int* A, int* B, int* C, int width) {
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    if (row < width && col < width) {
        C[row * width + col] = A[row * width + col] + B[row * width + col];
    }
}

int main() {
    // 矩陣維度
    int width = 4;

    // 分配CPU內(nèi)存
    int* A, * B, * C;
    A = (int*)malloc(width * width * sizeof(int));
    B = (int*)malloc(width * width * sizeof(int));
    C = (int*)malloc(width * width * sizeof(int));

    // 初始化A和B矩陣
    for (int i = 0; i < width; i++) {
        for (int j = 0; j < width; j++) {
            A[i * width + j] = i;
            B[i * width + j] = j;
        }
    }

    // 為GPU矩陣分配內(nèi)存
    int* d_A, * d_B, * d_C;
    cudaMalloc((void**)&d_A, width * width * sizeof(int));
    cudaMalloc((void**)&d_B, width * width * sizeof(int));
    cudaMalloc((void**)&d_C, width * width * sizeof(int));

    // 將矩陣從CPU內(nèi)存復(fù)制到GPU內(nèi)存
    cudaMemcpy(d_A, A, width * width * sizeof(int), cudaMemcpyHostToDevice);
    cudaMemcpy(d_B, B, width * width * sizeof(int), cudaMemcpyHostToDevice);

    // 配置CUDA核函數(shù)參數(shù)
    dim3 threads(width, width);
    dim3 grid(1, 1);
    matrixAdd10 <<<grid, threads >>> (d_A, d_B, d_C, width);

    // 等待CUDA核函數(shù)執(zhí)行完畢
    cudaDeviceSynchronize();

    // 將結(jié)果從GPU內(nèi)存復(fù)制到CPU內(nèi)存
    cudaMemcpy(C, d_C, width * width * sizeof(int), cudaMemcpyDeviceToHost);

    // 驗(yàn)證結(jié)果
    for (int i = 0; i < width; i++) {
        for (int j = 0; j < width; j++) {
            if (C[i * width + j] != i + j) {
                printf("錯(cuò)誤!");
                return 0;
            }
        }
    }
    printf("矩陣加法成功!");

    // 釋放CPU和GPU內(nèi)存
    free(A); free(B); free(C);
    cudaFree(d_A); cudaFree(d_B); cudaFree(d_C);
}

其實(shí)，CPU部分的main函數(shù)還是比較好懂的。核函數(shù)這邊就有點(diǎn)不知所措了，比如下面這兩行：

    int row = blockIdx.y * blockDim.y + threadIdx.y;
    int col = blockIdx.x * blockDim.x + threadIdx.x;

這兩行代碼用于計(jì)算當(dāng)前 CUDA 線程在二維矩陣中的位置。在 CUDA 編程模型中，我們通常將問(wèn)題劃分為多個(gè)線程塊 (block)，每個(gè)線程塊包含多個(gè)線程。線程塊和線程可以是一維、二維或三維的。在這個(gè)矩陣加法的例子中，我們使用二維線程塊和二維線程。

blockIdx 和 blockDim 分別表示線程塊索引和線程塊的尺寸，它們都是 dim3 類型的變量。threadIdx 表示線程的索引，也是一個(gè) dim3 類型的變量。x 和 y 分別表示這些變量的橫向和縱向分量。

int row = blockIdx.y * blockDim.y + threadIdx.y;

這行代碼計(jì)算當(dāng)前線程在二維矩陣中的行號(hào)。blockIdx.y 表示當(dāng)前線程所在的線程塊在縱向（行方向）上的索引，blockDim.y 表示每個(gè)線程塊在縱向上包含的線程數(shù)，threadIdx.y 表示當(dāng)前線程在所在線程塊中縱向的索引。將這些值組合在一起，可以計(jì)算出當(dāng)前線程在整個(gè)矩陣中的行號(hào)。

int col = blockIdx.x * blockDim.x + threadIdx.x;

這行代碼計(jì)算當(dāng)前線程在二維矩陣中的列號(hào)。blockIdx.x 表示當(dāng)前線程所在的線程塊在橫向（列方向）上的索引，blockDim.x 表示每個(gè)線程塊在橫向上包含的線程數(shù)，threadIdx.x 表示當(dāng)前線程在所在線程塊中橫向的索引。將這些值組合在一起，可以計(jì)算出當(dāng)前線程在整個(gè)矩陣中的列號(hào)。

通過(guò)這兩行代碼，我們可以為每個(gè)線程分配一個(gè)特定的矩陣元素，讓它執(zhí)行相應(yīng)的加法操作。這種并行計(jì)算方式可以顯著提高矩陣加法的計(jì)算速度。

這段代碼需要使用NVidia CUDA工具包中的nvcc來(lái)編譯了，我們將其存為matrix_add.cu:

nvcc -o matrix_add matrix_add.cu
./matrix_add

OpenCL

CUDA是一門NVidia專有的技術(shù)，在其它GPU上用不了。所以其它廠商一直在想辦法提供類似的技術(shù)。這其中，曾經(jīng)最被看好的就是OpenCL。OpenCL由Apple最初提出并由Khronos Group牽頭制定和管理標(biāo)準(zhǔn)。
OpenCL是一種用于編寫跨平臺(tái)的異構(gòu)計(jì)算程序的框架，支持使用C99, C++14和C++17語(yǔ)言編寫代碼，可以在多種處理器和操作系統(tǒng)上運(yùn)行，如CPU, GPU, DSP, FPGA等。
OpenCL的第一個(gè)版本于2008年發(fā)布。

我們來(lái)看下用OpenCL寫的計(jì)算矩陣加法的節(jié)選。
首先也是運(yùn)行在GPU上的核函數(shù)，然后通過(guò)enqueueNDRangeKernel將其放入執(zhí)行隊(duì)列中。

#include <iostream>
#include <vector>
#include <CL/cl.hpp>

const char* kernelSource = R"CLC(
__kernel void matrix_add(__global const int* A, __global const int* B, __global int* C, int rows, int cols) {
    int i = get_global_id(0);
    int j = get_global_id(1);

    int index = i * cols + j;

    if (i < rows && j < cols) {
        C[index] = A[index] + B[index];
    }
}
)CLC";

int main() {
    std::vector<std::vector<int>> A = {
        {1, 2, 3},
        {4, 5, 6},
        {7, 8, 9}
    };

    std::vector<std::vector<int>> B = {
        {9, 8, 7},
        {6, 5, 4},
        {3, 2, 1}
    };

    int rows = A.size();
    int cols = A[0].size();

    std::vector<int> A_flat(rows * cols), B_flat(rows * cols), C_flat(rows * cols);
    for (int i = 0; i < rows; ++i) {
        for (int j = 0; j < cols; ++j) {
            A_flat[i * cols + j] = A[i][j];
            B_flat[i * cols + j] = B[i][j];
        }
    }

    std::vector<cl::Platform> platforms;
    cl::Platform::get(&platforms);

    cl_context_properties properties[] = {
        CL_CONTEXT_PLATFORM, (cl_context_properties)(platforms[0])(), 0
    };
    cl::Context context(CL_DEVICE_TYPE_GPU, properties);

    cl::Program program(context, kernelSource, true);

    cl::CommandQueue queue(context);

    cl::Buffer buffer_A(context, CL_MEM_READ_ONLY, sizeof(int) * rows * cols);
    cl::Buffer buffer_B(context, CL_MEM_READ_ONLY, sizeof(int) * rows * cols);
    cl::Buffer buffer_C(context, CL_MEM_WRITE_ONLY, sizeof(int) * rows * cols);

    queue.enqueueWriteBuffer(buffer_A, CL_TRUE, 0, sizeof(int) * rows * cols, A_flat.data());
    queue.enqueueWriteBuffer(buffer_B, CL_TRUE, 0, sizeof(int) * rows * cols, B_flat.data());

    cl::Kernel kernel(program, "matrix_add");
    kernel.setArg(0, buffer_A);
    kernel.setArg(1, buffer_B);
    kernel.setArg(2, buffer_C);
    kernel.setArg(3, rows);
    kernel.setArg(4, cols);

    cl::NDRange global_size(rows, cols);
    queue.enqueueNDRangeKernel(kernel, cl::NullRange, global_size);

    queue.enqueueReadBuffer(buffer_C, CL_TRUE, 0, sizeof(int) * rows * cols, C_flat.data());

    std::vector<std::vector<int>> C(rows, std::vector<int>(cols));
    for (int i = 0; i < rows; ++i) {
        for (int j = 0; j < cols; ++j) {
            C[i][j] = C_flat[i * cols + j];
        }
    }

...

Direct3D

在Windows上，我們都知道微軟的主要用于游戲開(kāi)發(fā)的DirectX。
Direct X作為Windows直接訪問(wèn)硬件的游戲加速接口，早在1995年就推出了。不過(guò)Direct X 1.0的時(shí)候還不支持3D，只支持2D。因?yàn)榈谝粋€(gè)廣泛使用的3D加速卡3dfx Voodoo卡1996年才推出。
Direct3D 1.0于1996年問(wèn)世。不過(guò)這時(shí)候只是對(duì)標(biāo)OpenGL的框架，跟GPGPU關(guān)系還遠(yuǎn)著呢。

一直要到2009年，Windows 7時(shí)代的Direct3D 11.0，才正式可以支持計(jì)算著色器。Direct 3D 12.0于2015年和Windows 10同時(shí)代推出。

在Direct3D 12中，GPU指令是通過(guò)HLSL語(yǔ)言來(lái)寫的：

// MatrixAddition.hlsl

[numthreads(16, 16, 1)]
void main(uint3 dt : SV_DispatchThreadID, uint3 gt : SV_GroupThreadID, uint3 gi : SV_GroupID) {
    // 確保我們?cè)诰仃嚪秶鷥?nèi)
    if (dt.x >= 3 || dt.y >= 3) {
        return;
    }

    // 矩陣 A 和 B 的值
    float A[3][3] = {
        {1, 2, 3},
        {4, 5, 6},
        {7, 8, 9}
    };

    float B[3][3] = {
        {9, 8, 7},
        {6, 5, 4},
        {3, 2, 1}
    };

    // 計(jì)算矩陣加法
    float result = A[dt.y][dt.x] + B[dt.y][dt.x];

    // 將結(jié)果寫入輸出緩沖區(qū)
    RWStructuredBuffer<float> output;
    output[dt.y * 3 + dt.x] = result;
}

然后是CPU上的操作，要建立一個(gè)計(jì)算著色器，因?yàn)榧?xì)節(jié)比較多，我就略去了，只寫主干：

#include <d3d12.h>
#include <d3dcompiler.h>
#include <iostream>

// 創(chuàng)建一個(gè)簡(jiǎn)單的計(jì)算著色器的 PSO
ID3D12PipelineState* CreateMatrixAdditionPSO(ID3D12Device* device) {
    ID3DBlob* csBlob = nullptr;
    D3DCompileFromFile(L"MatrixAddition.hlsl", nullptr, nullptr, "main", "cs_5_0", 0, 0, &csBlob, nullptr);

    D3D12_COMPUTE_PIPELINE_STATE_DESC psoDesc = {};
    psoDesc.pRootSignature = rootSignature; // 假設(shè)已創(chuàng)建好根簽名
    psoDesc.CS = CD3DX12_SHADER_BYTECODE(csBlob);

    ID3D12PipelineState* pso = nullptr;
    device->CreateComputePipelineState(&psoDesc, IID_PPV_ARGS(&pso));

    csBlob->Release();
    return pso;
}

// 執(zhí)行矩陣加法計(jì)算
void RunMatrixAddition(ID3D12GraphicsCommandList* commandList, ID3D12Resource* outputBuffer) {
    commandList->SetPipelineState(matrixAdditionPSO);
    commandList->SetComputeRootSignature(rootSignature);
    commandList->SetComputeRootUnorderedAccessView(0, outputBuffer->GetGPUVirtualAddress());
    // 分發(fā)計(jì)算著色器，設(shè)置線程組的數(shù)量
    commandList->Dispatch(1, 1, 1);

    // 確保在繼續(xù)之前完成計(jì)算操作
    commandList->ResourceBarrier(1, &CD3DX12_RESOURCE_BARRIER::UAV(outputBuffer));
}

int main() {
    // 初始化 DirectX 12 設(shè)備、命令隊(duì)列、命令分配器等...
    // ...

    // 創(chuàng)建根簽名、PSO 和計(jì)算著色器相關(guān)資源
    // ...

    // 創(chuàng)建輸出緩沖區(qū)
    ID3D12Resource* outputBuffer = nullptr;
    device->CreateCommittedResource(
        &CD3DX12_HEAP_PROPERTIES(D3D12_HEAP_TYPE_DEFAULT),
        D3D12_HEAP_FLAG_NONE,
        &CD3DX12_RESOURCE_DESC::Buffer(3 * 3 * sizeof(float)),
        D3D12_RESOURCE_STATE_UNORDERED_ACCESS,
        nullptr,
        IID_PPV_ARGS(&outputBuffer)
    );

    // 創(chuàng)建并執(zhí)行命令列表
    ID3D12GraphicsCommandList* commandList = nullptr;
    device->CreateCommandList(0, D3D12_COMMAND_LIST_TYPE_DIRECT, commandAllocator, nullptr, IID_PPV_ARGS(&commandList));

    RunMatrixAddition(commandList, outputBuffer);

    // 關(guān)閉命令列表并執(zhí)行
    commandList->Close();
    ID3D12CommandList* commandLists[] = {commandList};
    commandQueue->ExecuteCommandLists(_countof(commandLists), commandLists);

    // 同步 GPU 和 CPU
    // ...

    // 從輸出緩沖區(qū)中讀取結(jié)果
    float result[3][3] = {};
    void* mappedData = nullptr;
    outputBuffer->Map(0, nullptr, &mappedData);
    memcpy(result, mappedData, sizeof(result));
    outputBuffer->Unmap(0, nullptr);

    // 輸出結(jié)果
    for (int i = 0; i < 3; ++i) {
        for (int j = 0; j < 3; ++j) {
            std::cout << result[i][j] << " ";
        }
        std::cout << std::endl;
    }

    // 清理資源
    // ...
}

Vulkan

Vulkan由Khronos Group牽頭制定和管理標(biāo)準(zhǔn)，是OpenGL的繼任者。它最早的技術(shù)來(lái)自于AMD。
Vulkan是一種用于編寫跨平臺(tái)的圖形和計(jì)算程序的框架，支持使用C和C++語(yǔ)言編寫代碼，可以在多種處理器和操作系統(tǒng)上運(yùn)行，如CPU, GPU, DSP, FPGA等。

Vulkan的1.0版本于2016年發(fā)布。

默認(rèn)情況下，Vulkan使用帶計(jì)算管線的glsl:

#version 450
#extension GL_ARB_separate_shader_objects : enable

layout (local_size_x = 16, local_size_y = 16, local_size_z = 1) in;

layout (binding = 0) readonly buffer InputA {
    float dataA[];
};

layout (binding = 1) readonly buffer InputB {
    float dataB[];
};

layout (binding = 2) writeonly buffer Output {
    float dataC[];
};

void main() {
    uint index = gl_GlobalInvocationID.x + gl_GlobalInvocationID.y * gl_NumWorkGroups.x * gl_WorkGroupSize.x;
    dataC[index] = dataA[index] + dataB[index];
}

然后，在主機(jī)程序中，完成以下步驟：

• 初始化Vulkan實(shí)例和物理/邏輯設(shè)備。
• 創(chuàng)建一個(gè)Vulkan計(jì)算管道，加載和編譯計(jì)算著色器。
• 為輸入矩陣A和B以及輸出矩陣C創(chuàng)建Vulkan緩沖區(qū)。
• 將輸入矩陣數(shù)據(jù)復(fù)制到輸入緩沖區(qū)。
• 創(chuàng)建描述符集布局和描述符池，以描述著色器中的資源綁定。
• 創(chuàng)建描述符集，并將輸入/輸出緩沖區(qū)綁定到描述符集中。
• 創(chuàng)建一個(gè)Vulkan命令緩沖區(qū)，以記錄計(jì)算著色器調(diào)度的命令。
• 開(kāi)始記錄命令緩沖區(qū)，并調(diào)用vkCmdBindPipeline和vkCmdBindDescriptorSets將計(jì)算管道和描- 述符集綁定到命令緩沖區(qū)。
• 使用vkCmdDispatch調(diào)度計(jì)算著色器執(zhí)行矩陣加法。
• 結(jié)束命令緩沖區(qū)記錄，將命令緩沖區(qū)提交到Vulkan隊(duì)列。
• 等待隊(duì)列執(zhí)行完成，并將輸出緩沖區(qū)的數(shù)據(jù)復(fù)制回主機(jī)內(nèi)存。
• 清理Vulkan資源。

具體代碼就不詳細(xì)列出了。
大致的代碼結(jié)構(gòu)為：

// Vulkan實(shí)例、設(shè)備、命令池、隊(duì)列
VkInstance instance;
VkDevice device;
VkCommandPool commandPool;
VkQueue queue;

// 矩陣維度
const int width = 4;

// 頂點(diǎn)緩沖區(qū)對(duì)象
VkBuffer vertexBuffer;
VkDeviceMemory vertexBufferMemory;

// 結(jié)果緩沖區(qū)對(duì)象
VkBuffer resultBuffer;
VkDeviceMemory resultBufferMemory;

// 著色器模塊和管線
VkShaderModule shaderModule;
VkPipeline pipeline;

// 創(chuàng)建頂點(diǎn)緩沖區(qū)
// 向緩沖區(qū)填充矩陣A和B
// ...

// 創(chuàng)建結(jié)果緩沖區(qū)
// 向緩沖區(qū)映射內(nèi)存 
void* resultData;
vkMapMemory(device, resultBufferMemory, 0, sizeof(int) * 4 * 4, 0, &resultData);

// 創(chuàng)建著色器模塊(矩陣加法著色器)  
const char* shaderCode = "上面的glsl"; 
shaderModule = createShaderModule(shaderCode);

// 創(chuàng)建圖形管線
// ...

// 記錄命令
VkCommandBuffer commandBuffer; 
VkCommandBufferAllocateInfo commandBufferAllocateInfo = ...;
vkAllocateCommandBuffers(commandPool, &commandBufferAllocateInfo, &commandBuffer);

// 開(kāi)始記錄命令
vkBeginCommandBuffer(commandBuffer, &beginInfo);

// 綁定頂點(diǎn)緩沖區(qū)和結(jié)果緩沖區(qū)
vkCmdBindVertexBuffers(commandBuffer, 0, 1, &vertexBuffer, &offset);
vkCmdBindBuffer(commandBuffer, 1, 0, resultBuffer, &offset);

// 繪制
vkCmdDraw(commandBuffer, 4, 1, 0, 0); 

// 結(jié)束記錄命令  
vkEndCommandBuffer(commandBuffer);

// 提交命令并執(zhí)行
VkSubmitInfo submitInfo = ...;
vkQueueSubmit(queue, 1, &submitInfo, VK_NULL_HANDLE);
vkQueueWaitIdle(queue); 

// 讀取結(jié)果矩陣
for (int i = 0; i < width; i++) {
    for (int j = 0; j < width; j++) {
        int result = ((int*)resultData)[i * width + j];
        printf("%d ", result);
    }
    printf("\n");
}

// 釋放Vulkan資源
...

WebGPU

WebGPU是剛剛要被Chrome瀏覽器支持的用于前端的GPU技術(shù)。
WebGPU是一種用于編寫跨平臺(tái)的圖形和計(jì)算程序的框架，支持使用JavaScript和WebAssembly語(yǔ)言編寫代碼，可以在多種瀏覽器和操作系統(tǒng)上運(yùn)行，如Chrome, Firefox, Safari等。WebGPU是由W3C的GPU for the Web工作組制定和管理標(biāo)準(zhǔn)，是WebGL的繼任者。
前面我們看到，源于NVidia技術(shù)的CUDA，源于Apple技術(shù)的OpenCL，源于微軟技術(shù)的DirectX，還有源于AMD技術(shù)的Vulkan在桌面和服務(wù)端百花爭(zhēng)艷。在移動(dòng)端自然也是少不了龍爭(zhēng)虎斗。

第一個(gè)提出WebGPU想法的是蘋果，2016年2月，蘋果公司提出了一個(gè)名為Web Metal的提案，旨在將Metal API的概念移植到Web平臺(tái)上。
2017年2月，微軟公司提出了一個(gè)名為Web D3D的提案，旨在將Direct3D 12 API的概念移植到Web平臺(tái)上。
2017年8月，Mozilla公司提出了一個(gè)名為Obsidian的提案，旨在創(chuàng)建一個(gè)基于Vulkan API的抽象層。

幾家爭(zhēng)執(zhí)不下，谷歌公司提出了一個(gè)名為NXT的提案，旨在創(chuàng)建一個(gè)基于Vulkan, Metal和Direct3D 12 API的抽象層。
2018年4月，W3C工作組決定將NXT作為規(guī)范草案的起點(diǎn)，并將其重命名為WebGPU。

既然是一個(gè)抽象層，著色器語(yǔ)言不管使用SPIR-V，Vulkan的GLSL，DirectX的HLSL或者蘋果的Metal Shading Language就都不合適了。
于是2019年，WebGPU社區(qū)組提出了一個(gè)新的著色器語(yǔ)言的提案，名為WebGPU Shading Language (WGSL)，旨在創(chuàng)建一個(gè)基于SPIR-V的文本格式，以提供一種安全、可移植、易于使用和易于實(shí)現(xiàn)的著色器語(yǔ)言。

下面的代碼展示下流程，這個(gè)時(shí)刻還有瀏覽器正式支持。等子彈飛一會(huì)兒瀏覽器正式上線了之后，我們?cè)诤竺鏁?huì)專門講。

看下圖：WebGPU的規(guī)范還沒(méi)release呢。WGSL的規(guī)范也同樣沒(méi)有最后release。

js
// 獲取WebGPU adapter和設(shè)備 
const adapter = await navigator.gpu.requestAdapter();
const device = await adapter.requestDevice();

// 矩陣維度
const width = 4;

// 創(chuàng)建緩沖區(qū) - 用作頂點(diǎn)緩沖區(qū)和結(jié)果緩沖區(qū)
const vertexBuffer = device.createBuffer({
  size: width * width * 4 * Int32Array.BYTES_PER_ELEMENT, 
  usage: GPUBufferUsage.VERTEX | GPUBufferUsage.STORAGE 
});

// 獲得緩沖區(qū)映射 - 填充矩陣A和B
const vertexBufferMapping = await vertexBuffer.map();
new Int32Array(vertexBufferMapping).fill(/* A和B矩陣 */);
vertexBuffer.unmap();

// 著色器代碼
const shaderCode = `
  kernel void addMatrices(device int* a [[buffer(0)]], 
                          device int* b [[buffer(1)]], 
                          device int* c [[buffer(2)]]) {
    const int width = 4;
    int tid = threadIdx.x * 4 + threadIdx.y;
    if (tid < width * width) {
      c[tid] = a[tid] + b[tid]; 
    }
  }
`;  

// 創(chuàng)建著色器模塊
const shaderModule = device.createShaderModule({ 
  code: shaderCode 
});

// 運(yùn)行著色器 - 執(zhí)行矩陣加法
const pipeline = device.createComputePipeline({
  compute: {
    module: shaderModule, 
    entryPoint: "addMatrices" 
  }
});
const passEncoder = device.createCommandEncoder();
const computePass = passEncoder.beginComputePass();   
computePass.setPipeline(pipeline);
computePass.setBuffer(0, vertexBuffer);  
computePass.setBuffer(1, vertexBuffer);
computePass.setBuffer(2, vertexBuffer);  
computePass.dispatch(1); 
computePass.endPass();
device.queue.submit([passEncoder.finish()]);

// 讀取結(jié)果 
const result = new Int32Array(
  await vertexBuffer.mapRead()
);
// 打印結(jié)果矩陣
... 

// 釋放資源

小結(jié)

雖然還沒(méi)有講細(xì)節(jié)，但是本篇為我們打開(kāi)了SIMD和GPU編程的一扇門。文章來(lái)源地址http://www.zghlxwxcb.cn/news/detail-432144.html

到了這里，關(guān)于2023年的深度學(xué)習(xí)入門指南(9) - SIMD和通用GPU編程的文章就介紹完了。如果您還想了解更多內(nèi)容，請(qǐng)?jiān)谟疑辖撬阉鱐OY模板網(wǎng)以前的文章或繼續(xù)瀏覽下面的相關(guān)文章，希望大家以后多多支持TOY模板網(wǎng)！