⚡

SYCL 実装の性能を比較してみた

Techblog管理者

2023/02/20に公開

GPU

sycl

tech

TL;DR: SYCL 2020 の実装を NVIDIA GPU で動かす手順の紹介と、畳み込み演算の性能比較
こんにちは。ソリューション第二事業部の坂部です。
以前（2017 年）当ブログにて、 SYCLを使ってOpenCLを単一ソースで書いてみる という記事が公開されました。 その記事は 2017 年時点の情報であり、2023 年の現時点ではいくぶんか状況が変わっていますので、当記事では新しい状況をお伝えします。

 SYCL とは単一の C++ ソースで、種々のアクセラレータ （GPU, FPGA など） を扱うための規格です。 SYCL を使うことで、ホスト向けとアクセラレータ向けのコードを別ファイルに分けることなく、 C++17 で記述できます。 また、CUDA や OpenCL と比較してレイヤの高いプログラミングパラダイムを持ち、ホスト・デバイス間のメモリ転送の記述を省略できる利点もあります。
規格の最新のバージョンは SYCL 2020 (Revision 6) です。
SYCL 2020 より過去のバージョン（SYCL 1.2 まで）は OpenCL に変換することを前提とした規格でしたが、 SYCL 2020 ではその他のバックエンドを実装できるようになりました。これにより、対応するアクセラレータの幅が広がっています。

 SYCL の実装Khronos の SYCL Overview サイト（2023/01/09 閲覧）によると、利用可能な実装は以下があります。
Intel oneAPI Data Parallel C++ (DPC++)
hipSYCL （注）
triSYCL
neoSYCL
Codeplay ComputeCpp
注: 2023 年 2 月に、 hipSYCL は Open SYCL に名称変更されました。 内容に変更はない ため、当記事では hipSYCL の名前のまま紹介します。
また、各実装がサポートするバックエンドやアクセラレータの一覧としては、 hipSYCL リポジトリのドキュメントにある図が分かりやすいです。 以下に引用します。


画像出典: https://github.com/illuhad/hipSYCL, Copyright (c) 2018 Aksel Alpay (BSD 2-Clause "Simplified" License)
本記事では、 NVIDIA GPU をターゲットにする場合に候補に挙がる
Intel oneAPI Data Parallel C++ (DPC++)
hipSYCL
の 2 通りの実装の使い方の紹介と、実装間の速度の比較を行いました。
NVIDIA GPU を使う場合において、 CUDA や OpenCL の環境構築に比べて SYCL の環境構築は複雑です。読者のみなさんが SYCL を試すときの手助けになれば幸いです。

 実験環境OS: Ubuntu 20.04.5 LTS
GPU: NVIDIA TITAN V (Volta)
CUDA 11.6

 Intel oneAPI Data Parallel C++ (DPC++) の環境構築Intel oneAPI Data Parallel C++ (DPC++) を NVIDIA GPU とともに使用する場合の手順を紹介します。
注: 2022年12月中旬、 oneAPI 2023.0 のリリースに際して oneAPI が NVIDIA GPU・AMD GPU のプラグインに対応したと発表されました。 当記事にこの内容は反映されておりません。

 コンパイラを入手し、ビルドするoneAPI DPC++ は apt などのパッケージシステムでインストールできますが、oneAPI 2023.0 以前は NVIDIA GPU の対応は GitHub で公開されている版のみでした。 GitHub 版 DPC++ を使用するには、まず DPC++ コンパイラをビルドする必要があります。
コンパイラをビルドする手順は Build DPC++ toolchain with support for NVIDIA CUDA に示されていますが、概要をまとめます。
「Prerequisites」 （Git, CMake, Python, Ninja, GCC） をインストールする
CUDA 11.6 をインストールする
作業ディレクトリを作り、リポジトリをクローンする
$ export DPCPP_HOME=~/sycl_workspace  
$ mkdir $DPCPP_HOME 
$ cd $DPCPP_HOME 
$ git clone https://github.com/intel/llvm -b sycl 
--cuda を指定してビルドする
$ python $DPCPP_HOME/llvm/buildbot/configure.py --cuda   
$ python $DPCPP_HOME/llvm/buildbot/compile.py
ビルドが完了し、 $DPCPP_HOME/llvm/build/bin/clang++ が生成されたことを確認する

 正しくデバイスを認識できるか確認する以下のサンプルコードを実行し、デバイス名とベンダー名を表示してみます。
使いたいデバイスの名前が表示されれば成功です。 DPC++ を使う環境ができました。
// show_names.cpp
#include <sycl/sycl.hpp>
#include <iostream>
int main() {
    // 使用可能デバイスの中からGPUを選ぶ
    auto device = sycl::device(sycl::gpu_selector_v);
    // デバイス名とベンダー名を取得し表示する
    std::cout << "Device Name: " << device.get_info<sycl::info::device::name>() << std::endl;
    std::cout << "Device Vendor: " << device.get_info<sycl::info::device::vendor>() << std::endl;
}
# コンパイル
$ $DPCPP_HOME/llvm/build/bin/clang++ show_names.cpp -fsycl -fsycl-targets=nvptx64-nvidia-cuda -o show_names
# 実行（動的リンクのため、ライブラリ探索パスに llvm/build/lib を加える）
$ LD_LIBRARY_PATH="$DPCPP_HOME/llvm/build/lib:${LD_LIBRARY_PATH:-}" ./show_names
Device Name: NVIDIA TITAN V
Device Vendor: NVIDIA Corporation

 hipSYCL の環境構築hipSYCL （Open SYCL）を NVIDIA GPU とともに使用する場合の手順を紹介します。

 LLVM/clang および hipSYCL コンパイラを入手し、ビルドする手順は Building and installing hipSYCL および hipSYCL installation instructions for CUDA に示されていますが、概要をまとめます。
「Software Dependencies」 （Python3, CMake, boost） をインストールする
CUDA 11.6 をインストールする
CUDA(NVPTX) サポートを有効にした LLVM/clang 14 をビルドする
LLVM/clang のビルド方法は Building LLVM with CMake を参照する
CUDA(NVPTX) サポート有効化のため、 configure 時に -D LLVM_TARGETS_TO_BUILD="X86;NVPTX" を付ける
hipSYCL リポジトリをクローンする
$ git clone https://github.com/illuhad/hipSYCL
$ cd hipSYCL
hipSYCL をビルドする
$ mkdir build
$ cd build
$ cmake -DCMAKE_INSTALL_PREFIX=$HOME/local -DCUDA_TOOLKIT_ROOT_DIR=/usr/local/cuda -DWITH_CUDA_BACKEND=ON -DBOOST_ROOT=/usr/include/boost ..
$ make install
$HOME/local/bin/syclcc がインストールされていることを確認する

 正しくデバイスを認識できるか確認する以下のサンプルコードを実行し、デバイス名とベンダー名を表示してみます。
使いたいデバイスの名前が表示されれば成功です。 hipSYCL を使う環境ができました。
// show_names.cpp （DPC++ の環境構築のものと同じコード）
#include <sycl/sycl.hpp>
#include <iostream>
int main() {
    // 使用可能デバイスの中からGPUを選ぶ
    auto device = sycl::device(sycl::gpu_selector_v);
    // デバイス名とベンダー名を取得し表示する
    std::cout << "Device Name: " << device.get_info<sycl::info::device::name>() << std::endl;
    std::cout << "Device Vendor: " << device.get_info<sycl::info::device::vendor>() << std::endl;
}
# LLVM/clang 14 をインストールしたディレクトリを syclcc 実行時に指定する
$ CLANG_ROOT=/opt/llvm-14/llvm-project/build/

# ビルドする
$ syclcc -O2 --hipsycl-targets='cuda:sm_70' --hipsycl-clang=$CLANG_ROOT/bin/clang++ --hipsycl-clang-include-path=$CLANG_ROOT/lib/clang/14.0.6 show_names.cpp -o show_names

$ ./show_names
Device Name: NVIDIA TITAN V
Device Vendor: NVIDIA

 性能比較この節では、簡単な画像処理を実装し、2 通りの実装 (oneAPI DPC++, hipSYCL) の速度を比較します。
画像処理の内容は、畳み込み処理とします。
SYCL を使って愚直に実装すると、次のようなコードになります。 SYCL の使い方の詳しい部分については当記事で扱いません。
#include <cassert>
#include <iostream>
#include <sycl/sycl.hpp>
#include <opencv2/core/core.hpp>
void convolution2d(const cv::Mat& src, cv::Mat& dst, const cv::Mat& kernel)
{
    using inputType = std::uint8_t;
    assert(src.type() == CV_8UC1);
    assert(src.type() == dst.type());

    using kernelType = float;
    assert(kernel.type() == CV_32FC1);

    const int halfSize = kernel.cols / 2;

    // 利用可能デバイスのうち GPU を選択する
    auto device = sycl::device(sycl::gpu_selector_v);
    // プロファイル（時刻の記録）を有効化したキューを生成する
    auto q = sycl::queue{device, sycl::property::queue::enable_profiling{}};

    // 諸画像の範囲とバッファオブジェクトを定義する
    const auto srcRows = src.rows;
    const auto srcCols = src.cols;
    const auto srcRange = sycl::range<2>{
        static_cast<std::size_t>(srcRows), static_cast<std::size_t>(srcCols)
    };
    auto bufferSrc = sycl::buffer<inputType, 2>{
        src.ptr<inputType>(), srcRange
    };
    auto bufferDst = sycl::buffer<inputType, 2>{
        dst.ptr<inputType>(),
        sycl::range<2>{static_cast<std::size_t>(dst.rows), static_cast<std::size_t>(dst.cols)}
    };
    auto bufferKernel = sycl::buffer<kernelType, 2>{
        kernel.ptr<kernelType>(),
        sycl::range<2>{static_cast<std::size_t>(kernel.rows), static_cast<std::size_t>(kernel.cols)}
    };

    // キューに関数を登録する
    auto ev = q.submit([&](sycl::handler& cgh){
        auto accessorSrc = bufferSrc.get_access<sycl::access_mode::read>(cgh);
        auto accessorDst = bufferDst.get_access<sycl::access_mode::write>(cgh);
        auto accessorKernel = bufferKernel.get_access<sycl::access_mode::read>(cgh);

        cgh.parallel_for(
            srcRange,
            [accessorSrc, accessorDst, accessorKernel, halfSize, srcRows, srcCols](sycl::item<2> it) {
                float sum = 0.0f;
                const auto y = it.get_id(0);
                const auto x = it.get_id(1);
                // 画像の端は計算をスキップする
                if (   y < halfSize || y >= (srcRows - halfSize)
                    || x < halfSize || x >= (srcCols - halfSize)) {
                    return;
                }
                // 畳み込みを計算する
                for (int dy = -halfSize; dy <= halfSize; dy++) {
                    for (int dx = -halfSize; dx <= halfSize; dx++) {
                        sum += accessorSrc[y + dy][x + dx] * accessorKernel[dy + halfSize][dx + halfSize];
                    }
                }
                // 結果を出力画像に書き込む
                accessorDst[y][x] = sum;
            }
        );
    });

    // 計算の開始時刻と終了時刻を取得・表示する
    const auto startTime =
        ev.get_profiling_info<sycl::info::event_profiling::command_start>();
    const auto endTime =
        ev.get_profiling_info<sycl::info::event_profiling::command_end>();
    std::cout << startTime << ";" << endTime << std::endl;
}

int main() {
    constexpr int width = 512;
    constexpr int height = 512;
    constexpr int kernelSize = 3;
    cv::Mat src{height, width, CV_8UC1};
    cv::Mat dst{height, width, CV_8UC1};
    cv::Mat kernel{kernelSize, kernelSize, CV_32FC1};
    // 画像とカーネルの中身をここで設定する（省略）
    convolution2d(src, dst, kernel);
}
畳み込みのように少々複雑な処理でも、 2 通りの実装で同じ結果を返す単一のコードを、容易に開発できています。
このコードでは SYCL のプロファイル機能を使って、計算の開始時刻と終了時刻を表示しています。
画像サイズとカーネルサイズをいくつか変えて実行した結果の表を以下に示します。 統計をとるための実行回数は 20 回とし、最初の 1 回に SYCL 処理系の初期化処理（DPC++ では 10 us 程度、 hipSYCL では 1 ms 程度）が含まれたため除外しています。 有効桁数を 2 桁で示しています。


Image Size
Kernel Size
DPC++ Median (us)
DPC++ Deviation (us)
DPC++ Throughput (GFLOPS)
hipSYCL Median (us)
hipSYCL Deviation (us)
hipSYCL Throughput (GFLOPS)


512x512
3x3
11
0.16
400
29
1.7
160

512x512
9x9
43
0.46
980
58
0.93
720

2048x2048
3x3
110
1.0
630
110
5.4
650

2048x2048
9x9
550
3.4
1200
540
1.4
1300

Throughput を見ると以下のことが分かります。
512x512 画像では、どちらのカーネルサイズでも DPC++ が速い （1.3 から 2.5 倍）
2048x2048 画像では、 DPC++ と hipSYCL は同等に速い

 まとめ今回の記事では以下を行いました。
Intel oneAPI DPC++ を使って NVIDIA GPU を扱う手順の紹介
hipSYCL を使って NVIDIA GPU を扱う手順の紹介
SYCL で畳み込みを実装し、同一 GPU で達成される速度を比較
問題サイズが小さいときには DPC++ が優位だったが、問題サイズを大きくすると大きな差がない

 参考文献SYCL 2020 API Reference Guide
SYCL™ 2020 Specification (revision 6)

Image Size	Kernel Size	DPC++ Median (us)	DPC++ Deviation (us)	DPC++ Throughput (GFLOPS)	hipSYCL Median (us)	hipSYCL Deviation (us)	hipSYCL Throughput (GFLOPS)
512x512	3x3	11	0.16	400	29	1.7	160
512x512	9x9	43	0.46	980	58	0.93	720
2048x2048	3x3	110	1.0	630	110	5.4	650
2048x2048	9x9	550	3.4	1200	540	1.4	1300

Fixstars Tech Blog /proc/cpuinfo

このブログは、高速化のエキスパートであるフィックスターズのエンジニアたちが、技術的な知見や検証結果を自由に発信しているテックブログです。「ソフトウェアの高速化・最適化」と「最新ハードウェア/AI技術の活用」に焦点を当てた、専門性の高い記事が中心です。

SYCL とは

SYCL の実装

実験環境

Intel oneAPI Data Parallel C++ (DPC++) の環境構築

コンパイラを入手し、ビルドする

正しくデバイスを認識できるか確認する

hipSYCL の環境構築

LLVM/clang および hipSYCL コンパイラを入手し、ビルドする

正しくデバイスを認識できるか確認する

性能比較

まとめ

参考文献

Discussion