769 views

この記事は最終更新から 389日 が経過しています。

Single processor, Multi processorではありません。
Single-core, Multi-coreでの比較です。

【1】やりたいこと

過去記事 (46)【Othello AI】オセロAIにも相性がある。 では、CPU実行していたプログラムを RTX 5070ti に実行させたところ、処理速度が 10倍近く向上した。

あらためて GPUは凄いと驚いたが、
CPUとGPUの差がどれほどなのか？
の実感が持てない。

そこで、今回はシンプルな 1次元配列（要素数10⁹個）の足し算で、速度性能を比較してみることにした。

【2】やってみた

1) 比較する環境

以下の 3環境で同じプログラムを実行し、その結果を比較する。

#	プロセッサ	備考
1	CPU	single thread実行
2	CPU	multi thread実行（OpenMP）
3	GPU	ブロック数指定は最大値である 1024

因みに、使用する CPU, GPUは以下の通り。

CPU	Intel Core Ultle 7 265KF (20cores)
GPU	NVIDIA GeForce RTX 5070ti (8960cores)

2) 実験内容

4バイト浮動小数点数の配列の足し算をやらせる。
形式言語的に書けばこんな感じだ。

float A[1_000_000_000], B[1_000_000_000], C[1_000_000_000];
C[i] = A[i] + B[i];  // for i = 0 to 999_999_999

もしくは、こんな感じだ。

∀ i ∈ [0, N):  C[i] := A[i] + B[i]

配列の要素数は 10⁹個、すなわち 10億個 だ。

「4バイト浮動小数点数の計算を手計算で 10億回やってください。10億円あげます。」
と言われても絶対に断る。

人力で 5秒に一つ、
寝ずに計算したとして 158年かかる・・・

3) 実験用プログラム

CPU用、GPU用の二種類のプログラムを用意した。
CPU版の single-core or multi-core は、コンパイルスイッチ指定で分ける。

(1) CPU用実験コード

注意：
このプログラムでは、sizeof(float) x 10億個 x 3セット = 12GB のメモリを allocateしている。
メモリ搭載量の小さいシステムで実行する場合は、個数を減らす必要がある。

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

#define N 1000000000

#ifdef USE_OPENMP
#include <omp.h>
#endif

/*//////////////////////////////////////////////////////////////////////////////*/
int main() {
    size_t size = N * sizeof(float);
    /*--------------------------------------------------------------------------*/
    /* データ作成 */
    float *A = malloc(size);
    float *B = malloc(size);
    float *C = malloc(size);
    for (int i = 0; i < N; i++) {
        A[i] = 1.0f;
        B[i] = 2.0f;
    }
    /*--------------------------------------------------------------------------*/
    struct timespec start, end;
    clock_gettime(CLOCK_MONOTONIC, &start);     /* 時間計測開始 */
    /****************************************************************************>>> 計測区間 */
    /* 演算実行 */
    #pragma omp parallel for
    for (int i = 0; i < N; i++) {
        C[i] = A[i] + B[i];
    }
    /****************************************************************************<<< 計測区間 */
    clock_gettime(CLOCK_MONOTONIC, &end);       /* 時間計測終了 */
    long sec_diff  = end.tv_sec  - start.tv_sec;
    long nsec_diff = end.tv_nsec - start.tv_nsec;
    double elapsed_ms = sec_diff * 1000.0 + nsec_diff / 1e6;
    printf("time: %.3f ms\n", elapsed_ms);
    /*--------------------------------------------------------------------------*/
    /* メモリ解放 */
    free(A); free(B); free(C);
    /*--------------------------------------------------------------------------*/
    return 0;
}

CPU用プログラムをコンパイルし、Single thread用プログラム cpu_1 を生成する。

$ gcc cpu.c -o cpu_1

CPU用プログラムをコンパイルし、Multi thread用プログラム cpu_N を生成する。
#pragma omp parallel for と書いたところがマルチスレッド実行コードに展開される。

$ gcc cpu.c -o cpu_N -fopenmp

(2) GPU用実験コード

CPU版と比較すると、以下の処理が増えた分だけコード量が多い。
・GPUメモリ確保
・CPU → GPUデータ転送（INPUT）
・GPU用演算コード (CUDA)
・GPU → CPUデータ転送（OUTPUT）

まずは、上記のオーバーヘッド部分を無視し、
純粋にベクトル計算だけを対象に速度計測してみる。

注意：
このプログラムでは、ホスト(CPU)側、デバイス(GPU)側でそれぞれに
sizeof(float) x 10億個 x 3セット = 12GB のメモリを allocateしている。
メモリ搭載量の小さいシステムで実行する場合は、個数を減らす必要がある。

#include <stdio.h>
#include <cuda_runtime.h>

#define N 1000000000
#define THREADS_PER_BLOCK 1024

/*//////////////////////////////////////////////////////////////////////////////*/
__global__ void vector_add(float *A, float *B, float *C, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) {
        C[i] = A[i] + B[i];
    }
}

/*//////////////////////////////////////////////////////////////////////////////*/
int main() {
    size_t size = N * sizeof(float);
    /*--------------------------------------------------------------------------*/
    /* データ作成 */
    float *A = (float*)malloc(size);
    float *B = (float*)malloc(size);
    float *C = (float*)malloc(size);
    for (int i = 0; i < N; i++) {
        A[i] = 1.0f;
        B[i] = 2.0f;
    }
    /*--------------------------------------------------------------------------*/
    /* CPU → GPUデータ転送 */
    float *d_A, *d_B, *d_C;
    cudaMalloc(&d_A, size);
    cudaMalloc(&d_B, size);
    cudaMalloc(&d_C, size);
    cudaMemcpy(d_A, A, size, cudaMemcpyHostToDevice);
    cudaMemcpy(d_B, B, size, cudaMemcpyHostToDevice);
    /*--------------------------------------------------------------------------*/
    struct timespec start, end;
    clock_gettime(CLOCK_MONOTONIC, &start);     /* 時間計測開始 */
    /****************************************************************************>>> 計測区間 */
    /* 演算実行 */
    int blocks = (N + THREADS_PER_BLOCK - 1) / THREADS_PER_BLOCK;
    vector_add <<< blocks, THREADS_PER_BLOCK >>> (d_A, d_B, d_C, N);
    cudaDeviceSynchronize();     /* GPU側の処理完了と同期 */
    /****************************************************************************<<< 計測区間 */
    clock_gettime(CLOCK_MONOTONIC, &end);       /* 時間計測終了 */
    long sec_diff  = end.tv_sec  - start.tv_sec;
    long nsec_diff = end.tv_nsec - start.tv_nsec;
    double elapsed_ms = sec_diff * 1000.0 + nsec_diff / 1e6;
    printf("time: %.3f ms\n", elapsed_ms);
    /*--------------------------------------------------------------------------*/
    /* GPU → CPUデータ転送 */
    cudaMemcpy(C, d_C, size, cudaMemcpyDeviceToHost);
    /*--------------------------------------------------------------------------*/
    /* メモリ解放 */
    cudaFree(d_A); cudaFree(d_B); cudaFree(d_C);
    free(A); free(B); free(C);
    /*--------------------------------------------------------------------------*/
    return 0;
}

GPU用ソースファイルは nvcc（NVIDIA Cuda Compiler） でコンパイルし、実行プログラム gpu を生成する。

$ nvcc gpu.cu -o gpu

4) 実行結果

前述の 3種類のプログラムの実行結果は以下の通り。

$ ./cpu_1
time: 1536.077 ms

$ ./cpu_N
time: 206.356 ms

$ ./gpu
time: 15.783 ms

個人的な期待値よりはだいぶ低かったが GPUが高速だ。

GPUは、CPU（single-core）の 97倍
GPUは、CPU（20 multi-cores）の 13倍

5) データ転送のオーバーヘッド時間を含めて計測してみる。

上の方で、
まずは、上記のオーバーヘッド部分を無視し、純粋にベクトル計算だけを対象に速度計測してみる。
と書いたが、実運用する場合には、このオーバーヘッド部分を無視するわけにはいかない。

処理時間の計測範囲をメモリ転送を含む範囲に変更してみた。
このプログラムを実行すると・・・

$ ./gpu
time: 2020.624 ms

ん？？？

CPU（Single thread）よりも遅くなったぞ・・・

以下の個所に +2005[ms] の時間がかかったということだ。

    /* CPU → GPUデータ転送 */
    float *d_A, *d_B, *d_C;
    cudaMalloc(&d_A, size);
    cudaMalloc(&d_B, size);
    cudaMalloc(&d_C, size);
    cudaMemcpy(d_A, A, size, cudaMemcpyHostToDevice);
    cudaMemcpy(d_B, B, size, cudaMemcpyHostToDevice);
：
    /* GPU → CPUデータ転送 */
    cudaMemcpy(C, d_C, size, cudaMemcpyDeviceToHost);

メインのベクトル計算自体はたったの 15.783[ms] なのに・・・
おそるべしメモリ転送コスト・・・

【3】所感

CUDAでプログラムを実装する際には、
CPU(Host)-GPU(Device)間のデータ転送コストに注意しましょう。
CPU(Host)-GPU(Device)間のデータ転送の量と頻度を減らせるように、設計を最適化しましょう。
ということだな。

続きはこちら・・・
(49) Python vs C言語でCUDAベクトル計算速度を比較する。

---