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

// 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

// 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

#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);
}