ITエンジニアの技術メモ: 並列処理

背景

OpenMPを利用して並列化を行い、処理を高速化する必要が出てきたため、OpenMPで並列化処理を記述する方法について記述する

記事の目的

OpenMPで並列化処理を行う際のテンプレートを作成する

OpenMP

ここでは、OpenMPを利用したCプログラムの記述方法について記載する。

OpenMPとは

OpenMPは、OpenMP ARBが提供する並列コンピューティング環境において共有メモリのマルチスレッド処理をサポートするために開発されたAPIである。

利点

マルチスレッド並列なプログラムをディレクティブを挿入するだけで実現できる
gccなどで標準サポートされており、導入が容易である

テンプレート

OpenMPを利用して並列処理を行うテンプレートを記載する。

// CMakeLists.txt　に追記
# For OpenMP
find_package(OpenMP REQUIRED)
if(OpenMP_FOUND)
    message(STATUS "Use OpenMP")
    set(CMAKE_C_FLAGS "${CMAKE_C_FLAGS} ${OpenMP_C_FLAGS}")
    set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} ${OpenMP_CXX_FLAGS}")
endif()

// openmp_sample_code.cpp
#include <iostream>
#ifdef _OPENMP
// ヘッダーファイルをインクルード
#include <omp.h>
#endif
int main()
{
#ifdef _OPENMP
    // スレッド数を設定(今回は4スレッド)
    omp_set_num_threads(4);
#endif
#ifdef _OPENMP
    // スレッド数取得
    std::cout << "The number of processors is "<< omp_get_num_procs() << std::endl;
#endif
#ifdef _OPENMP
    // ここから並列処理を行う
    // 【注】カッコ内の処理は、スレッド数回実行される
    #pragma omp parallel
#endif
    {
        int a = 0, b= 0;
#ifdef _OPENMP
        // section毎に並列処理を行う
        #pragma omp sections
#endif
        {
#ifdef _OPENMP
            #pragma omp section
#endif
            {
                // 並列化される処理 1
                a ++; 
            }
#ifdef _OPENMP
            #pragma omp section
#endif 
            {
                // 並列化される処理 2
                b ++; 
            }
        }
#ifdef _OPENMP
        // ここまでの並列処理が完了するまで待機
        #pragma omp barrier
        // シングルスレッド処理
        // 【注】マルチスレッド処理内におけるプライベート変数は引き継がれない場合がある
        //  このコードでは、aまたはbのどちらかが0になる
        #pragma omp single
#endif
        {
            std::cout << "a=" << a << std::endl;
            std::cout << "b=" << b << std::endl;
        }
    }
    size_t i;
    int c = 0;
    int sum = 0;
#ifdef _OPENMP
    // forループを並列化
    // 【注】 #pragma omp parallel内で#pragma omp parallel forすると、2重に並列化するため注意
    // 　　　 上記の場合、#pragma omp parallel forを#pragma omp forとする
    // iを各スレッドのプライベート変数に指定
    // 並列化した際にスレッドをを静的(static)に割り当て、チャンクサイズを2に指定
    //        (この場合、i=0,1,4,5,...とi=2,3,6,7,...でスレッドが作られる)
    //        (動的に割り当てる場合はdynamicとする。ただし、スレッド生成コストが増大する)
    // sumは、i*iの処理が並列で行われた後、+=の処理がまとめて行われる
    // orderedは、forループ内にorderedセクションがあることを示す
    #pragma omp parallel for private(i) schedule(static, 2) reduction(+, sum) ordered
#endif 
    for(i = 0; i < 10; ++i)
    {
       sum += i*i;
#ifdef _OPENMP
        // criticalセクション内は、最大で１つのスレッドしか実行されない
        // ただし、スレッド待ちの処理コストが発生する
        #pragma omp critical
#endif
        {
            c += i*i;
            std::cout << "c=" << c << std::endl;
        }
#ifdef _OPENMP
        // orderedセクション内は、順番に実行される
        #pragma omp ordered
#endif
        {
            std::cout << "i=" << i << std::endl;
        }
        std::cout << "sum=" << sum << std::endl;
    }
}

まとめ

OpenMPでプログラミングをする際のテンプレートを調査、記載した

参考文献

変更履歴

2019/10/23: 新規作成
2019/12/15: コード修正

背景

CUDAを利用して並列化を行い、処理を高速化する必要が出てきたため、CUDAで並列化処理を記述する方法について記述する

記事の目的

CUDAで並列化処理を行う際のテンプレートを作成する

CUDA

ここでは、CUDAを利用したCプログラムの記述方法について記載する。

CUDAとは

CUDAは、nvidia社が提供するGPUを利用した並列演算プログラミング基盤である

利点

ユーザーが多く、情報も入手しやすい
ライブラリが充実している
導入が容易である

テンプレート

CUDAを利用して並列処理を行うテンプレートを記載する。

// cuda_sample_code.cu
#include <stdio.h>

// デバイス(GPU)側のスレッドの設定 //////////////////////////////////
// BLOCK_SIZE_Xは、1ブロックあたりのスレッド数を表す
// GPUの種類により、1ブロックあたりのスレッド数の制限が異なる
// 最適なスレッド数を設定しないと、カーネル関数の処理がスキップされる
// 注)上記の場合、エラーでプロセスが落ちる事はない
#define BLOCK_SIZE_X 512

// カーネル関数(GPUで処理する関数)vec_sumの宣言 /////////////////////
// この関数がGPUで並列実行される
__global__
void vec_sum(float k, float *a, float *b, float *c)
{
    // iは、"for(int i=0; i < grid.x * block.x; i++)" の値を取る
    // ただし、並列処理のため、i++順に処理されるわけではない
    int i = blockIdx.x*blockDim.x + threadIdx.x;
    c[i] = k*a[i] + b[i];

    // GPU内のスレッドの同期 ////////////////////////////////////////
    // block_size * grid_size個の全てのスレッドが、ここで同期する
    __syncthreads();
}

// カーネル関数を呼び出す関数を宣言 /////////////////////////////////
void cuda_handler(float *a, float *b, float *c)
{
    // デバイス用のポインタを宣言 ///////////////////////////////////
    float *d_a, *d_b, *d_c;

    // デバイス(GPU)側の配列のメモリを確保 //////////////////////////
    // デバイス側のメモリ確保処理は重いため、回数は減らした方が良い
    cudaMalloc(&d_a, N*sizeof(float));
    cudaMalloc(&d_b, N*sizeof(float));
    cudaMalloc(&d_c, N*sizeof(float));

    // ホスト側の配列内のデータをデバイス側にコピー /////////////////
    cudaMemcpy(d_a, a, N*sizeof(float), cudaMemcpyHostToDevice);
    cudaMemcpy(d_b, b, N*sizeof(float), cudaMemcpyHostToDevice);

    // デバイス側の配列cの全ての要素に0をセット /////////////////////
    cudaMemset(d_c, 0, N*sizeof(float));

    // cudaを利用した処理のうち、最後のエラーを取得し、表示 /////////
    // ここでは、メモリへのデータセットのエラーチェックに利用している
    // 注)nvccのコンパイルオプションに"-g -G"を追加しないと動作しない
    // 注)エラーチェックをアクティブにすると、性能が極端に落ちる
    checkCudaErrors(cudaGetLastError());

    // 並列処理するスレッド数を定義 /////////////////////////////////
    // 総スレッド数は block_size * grid_size = N 個である
    // x, y, zの3次元まで設定可能である
    dim3 block_size (BLOCK_SIZE_X, 1, 1);
    dim3 grid_size  (N / block.x, 1, 1);

    // カーネル関数(GPUで処理する関数)の呼び出し ////////////////////
    // カーネル関数内では、デバイス側のメモリ内のデータのみ操作可能
    vec_sum<<<grid_size, block_size>>>(2.0f, d_a, d_b, d_c);

    // cudaを利用した処理のうち、最後のエラーを取得し、表示 /////////
    // ここでは、vec_sumのエラーチェックに利用している
    // 注)nvccのコンパイルオプションに"-g -G"を追加しないと動作しない
    // 注)エラーチェックをアクティブにすると、性能が極端に落ちる
    checkCudaErrors(cudaGetLastError());

    // この行までに実行されたカーネル関数の処理が完了するまで待機 ///
    // デバイスとホストの処理は非同期である
    // 同期処理を行うか、cudaMemcpyするまで、互いは独立して動作する
    cudaThreadSynchronize();

    // 計算結果をホスト側にコピー ///////////////////////////////////
    cudaMemcpy(c, d_c, N*sizeof(float), cudaMemcpyDeviceToHost);

    // デバイス(GPU)側の配列のメモリを開放 //////////////////////////
    cudaFree(d_a);
    cudaFree(d_b);
    cudaFree(d_c);

    // cudaを利用した処理のうち、最後のエラーを取得し、表示 /////////
    // ここでは、メモリへのデータセットのエラーチェックに利用している
    // 注)nvccのコンパイルオプションに"-g -G"を追加しないと動作しない
    // 注)エラーチェックをアクティブにすると、性能が極端に落ちる
    checkCudaErrors(cudaGetLastError());
}

int main(void)
{
    // 計算回数の設定 ///////////////////////////////////////////////
    // N = 512×2048
    int N = 1<<20;

    // ホスト用のポインタを宣言 /////////////////////////////////////
    float *a, *b, *c;

    // ホスト側の配列のメモリを確保 /////////////////////////////////
    a = (float*)malloc(N*sizeof(float));
    b = (float*)malloc(N*sizeof(float));
    c = (float*)malloc(N*sizeof(float));

    // a, bの配列にそれぞれ1，2を代入////////////////////////////////
    for (int i = 0; i < N; i++) {
        a[i] = 1.0f;
        b[i] = 2.0f;
    }

    // cudaでの処理を行う関数 ///////////////////////////////////////
    cuda_handler(a, b, c);

    // 計算結果の確認 ///////////////////////////////////////////////
    float maxError = 0.0f;
    for (int i = 0; i < N; i++)
    {
        maxError = max(maxError, abs(c[i]-4.0f));
    }
    printf("Max error: %f", maxError);

    // ホスト側の配列のメモリを開放 /////////////////////////////////
    free(a);
    free(b);
    free(c);

    return 0;
}

テンプレートのコンパイル方法

テンプレートのコンパイル方法を記載する。

nvcc -lcuda ./cuda_sample_code.cu -o ./cuda_sample

デバッグログを出力する場合は、下記のように記載する。

nvcc -lcuda -g -G ./cuda_sample_code.cu -o ./cuda_sample

テンプレートの実行結果

テンプレートの実行結果を記載する。

$ ./cuda_sample
Max error: 0.000000

まとめ

CUDAでプログラミングをする際のテンプレートを調査、記載した

参考文献

変更履歴

2019/08/26: 新規作成

ITエンジニアの技術メモ

2019/10/23

OpenMPを使用した並列化

背景

記事の目的

OpenMP

OpenMPとは

利点

テンプレート

まとめ

参考文献

変更履歴

2019/08/26

forループをCUDAで並列化する方法

背景

記事の目的

CUDA

CUDAとは

利点

テンプレート

テンプレートのコンパイル方法

テンプレートの実行結果

まとめ

参考文献

変更履歴

MQTTの導入

ページビュー