SYCL プログラミング入門：初心者向けガイド

2025年4月3日

はじめに

ソリューション第二事業部の武本です。

当ブログにて、SYCLを使ってOpenCLを単一ソースで書いてみる (2017年), SYCL 実装の性能を比較してみた (2023年)という記事が以前公開されました。 これらの記事ではSYCL環境構築方法やSYCLのソース等が記載されましたが、その前の段階の SYCL の概要・使用方法についてはあまりなされていなかったので、当記事では SYCL について概要の説明から最低限理解必要な基本機能と用語、簡単なプログラムの作成についてを説明します。当記事は、SYCL2020 準拠で作成しました。

対象読者

  • SYCL に興味がある方
  • 基礎的なC++を理解している方
  • 基礎的な並列プログラミングを理解している方
    • CUDA との比較内容もあるので、CUDA の知識があればより理解は容易となります

SYCL とは

公式には

SYCL は異種システムをプログラミングするオープンな業界標準です。SYCL の設計により、標準の C++ ソースコードを異種デバイスまたはホストのいずれかで実行できるように記述できます。

とあります。すなわち、SYCL は様々なデバイスを扱うことができるように用意された統一規格です。

本来、 CPU であったり、NVIDIA 製 GPU 、AMD 製 GPU 等に同じ処理をさせたい場合はそれぞれ異なる言語とフレームワークを用いて開発する必要がありますが、 SYCL でコードを記述することでこれらのすべてのデバイスで動作させることができます。

SYCL を用いる代表的なメリットに以下のようなものがあります。

  • これ一つで様々なデバイスでの処理を記述できる
  • 単一のソースコードで記述可能
    • デバイスで実行してほしいコードも .cpp ファイル単体に記述することができる
  • SYCL2020 においては C++17 を使用可能

簡単なプログラム

ごく簡単な動作をするプログラムを記載し、このプログラムをもとにして SYCL の基本的な構文や機能を説明します。

#include <sycl/sycl.hpp>
#include <iostream>

int main() {
    constexpr int N = 10;
    int arr[N] = {0};

    {
        // バッファーはホストとデバイス間でデータ転送をする方法の１つ
        sycl::buffer<int, 1> buf(arr, sycl::range<1>(N));

        // コマンドグローバルキューを取得する
        sycl::queue myQueue;

        // デバイス上に変数を展開する
        myQueue.submit([&](sycl::handler &cgh) {
            // アクセサを取得する
            sycl::accessor acc(buf, cgh, sycl::write_only);

            // デバイス上で実行する処理を記述する
            cgh.parallel_for(sycl::range<1>(N), [=](sycl::item<1> id) {
                // デバイス上で実行
                acc[id.get_linear_id()] = id.get_linear_id();
            });
        })// このタイミングで arr に buf の内容がコピーされる
    }

    for (int i = 0; i < N; i++) {
        std::cout << "stdout_result: " << arr[i] << "\n" << std::endl;
    }

    return 0;
}

出力結果

stdout_result: 0
stdout_result: 1
stdout_result: 2
stdout_result: 3
stdout_result: 4
stdout_result: 5
stdout_result: 6
stdout_result: 7
stdout_result: 8
stdout_result: 9

上記のプログラムは、単純にCPUで処理するコードで記載すると以下のようになります。

#include <iostream>

int main() {
    constexpr int N = 10;
    int arr[N] = {0};

    for (int i = 0; i < N; i++) {
        arr[i] = i;
    }

    for (int i = 0; i < N; i++) {
        std::cout << "stdout_result: " << arr[i] << "\n" << std::endl;
    }

    return 0;
}

ホスト (host) とデバイス (device)

SYCL プラットフォーム・モデルは、OpenCL プラットフォーム・モデルをベースにしています。モデルは、デバイスと呼ばれる 1 つ以上の異種デバイスに接続されたホストで構成されます。

公式ドキュメントではこのように記載されています。 上記のプログラムのコメントで多用された用語ですが、ホスト/デバイスとは以下のことを指します。

  • ホスト：制御用のプロセッサ
  • デバイス：演算用プロセッサ

主にホストは CPU 、デバイスは CPU, GPU, DSP などとなります。 例えば、CUDA ではホストが CPU 、デバイスが GPU となります。

後述しますが、 cgh.parallel_for の第 2 引数に指定された無名関数がデバイス上で実行されます

また、 OpenCL プラットフォームと記載されていますが、 OpenCL とは、複数のデバイスを制御するための API です。詳細は公式サイトを参照してください。

SYCL カーネル関数（SYCL kernel function）

デバイスで実行できる関数オブジェクトは、SYCL カーネル関数と呼ばれます。

公式に記載の通り、デバイスで実行するように記載した関数のことを SYCL カーネル関数と呼びます。 CUDA で言えば、 __global__ 属性が付与された関数に相当します。

上記のプログラムの cgh.parallel_for の第 2 引数に指定された無名関数が SYCL カーネル関数です。

コマンドグループ (command group)

SYCL では、デバイス上のデータを処理する操作は、コマンドグループ関数オブジェクトを使用して表現されます。各コマンドグループ 関数オブジェクトには、カーネルを使用してデバイス上でデータを正しく処理するのに必要なすべてのワークを実行するコマンド・グ ループ・ハンドラーがあります。この方法では、データを転送および処理するコマンドグループは、実行するデバイス上のコマンドグ ループとしてエンキューされます。コマンドグループは、自動的に SYCL キューに送信されます。

とあり、簡潔に示すとコマンドグループとは「デバイス上のデータを処理する操作」のことを指します。

上記のプログラムで submit の引数に指定された無名関数がコマンドグループです。 SYCL カーネル関数に、アクセサ等を加えたより広い範囲の関数オブジェクトであるといえます。

sycl::queue

SYCL コマンドキューは、SYCL デバイスで実行されるコマンドグループを保持するオブジェクトです。

とあり、コマンドグループを保持しておくデータ構造（オブジェクト）です。 sycl::queue は以下のように「どのデバイス上で実行するコマンドグループを保持するのか」の情報を付与して宣言されることが多いです。

// 1. デフォルトのデバイス（GPU, CPU or ...）で実行するコマンドグループを保持
sycl::queue q(sycl::default_selector_v);
// 2. 1 と同じ意味
sycl::queue q;
// 3. デフォルトの CPU で実行するコマンドグループを保持
sycl::queue q(sycl::cpu_selector_v);
// 4. デフォルトの GPU で実行するコマンドグループを保持
sycl::queue q(sycl::gpu_selector_v);

submit

SYCL queue を使用して submit メンバー関数送信により SYCL ランタイムで実行されるコマンドグループを送信できます。

とあるように、コマンドグループの関数オブジェクトを sycl::queue に格納（push）する際に用います。

カーネルを実行するデバイスの確認

以下のようなコードを記述することで、どのデバイスで、どのプラットフォームで実行しているかをプログラム実行時に確認することができます。

sycl::queue q;
std::cout << q.get_device().get_info<sycl::info::device::name>() << std::endl; // デバイス名
std::cout << q.get_device().get_info<sycl::info::device::vendor>() << std::endl; // デバイスベンダー
std::cout << q.get_device().get_platform().get_info<sycl::info::platform::name>() << std::endl; // プラットフォーム名
std::cout << q.get_device().get_platform().get_info<sycl::info::platform::vendor>() << std::endl; // プラットフォームベンダー

sycl::buffer

buffer クラスは、SYCL カーネルで使用されるアクセサークラスを介してアクセスする必要がある 1 次元、2 次元、または 3 次元の共有配列を定義します。

SYCL カーネル関数内で利用したい配列（1, 2, 3 次元まで可）は sycl::buffer として定義することができます。 ただし、sycl::buffer として定義するだけでは利用できず、以下で説明するアクセサを介してのみ利用することができます。

上記のサンプルプログラムではホストの配列 arr のためのバッファとして buf という sycl::buffer を定義しています。

アクセサ

デバイスから様々なメモリにアクセスするための機能をアクセサと呼びます。 「様々なメモリ」と記載しましたが、バッファ(sycl::buffer)のほかに、ホストのみ、あるいはデバイスのみがアクセスできるメモリがあります。 デバイスのみがアクセスできるメモリの例としては、「ローカルメモリ」といったものがあり、これらのメモリにアクセスするための機能です。

sycl::accessor

sycl::accessor は、アクセサの一種で、 sycl::buffer にアクセスするためのクラスです。

上記のプログラムでは、sycl::buffer である buf に読み取り専用でアクセスするためのアクセサとして acc という sycl::accessor を定義しています。 ここで、読み取り専用と書きましたが、実際には sycl::accessor のコンストラクタの第 3 引数に以下のように指定することで読み込み、書き込み処理の可否を指定することができます。

sycl::read_only // 読み込み専用
sycl::write_only // 書き込み専用
sycl::read_write // 読み書き可能

このようにアクセサを介して sycl::buffer の中身を変更しますが、この変更はホスト側の配列(上記のプログラムでは arr)の配列には即座には反映されません。cuda では cudaMemcpy により明示的にデバイスからホストへデータをコピーしますが、上記のプログラムの sycl で反映されるタイミングは sycl::accessor のデストラクタが呼ばれた際（生存期間が切れた際）となります。この時にデストラクタの動作により自動的に数値がホスト側の配列にコピーされます。

parallel_for（データ並列カーネル）

ここまででデバイス上での配列(データ)へのアクセス方法について分かりました。ここで、実際にこのデータに処理を加える(デバイス上で処理をする)ための関数の一つが parallel_for です。 parallel_for によって SYCL カーネル関数を呼び出すことができ、パラメータを指定して並列処理を行うことができます。

並列処理を行うにあたり、処理単位を指定し、各処理単位ではどのデータを処理するかを指定する必要があります。SYCL ではこれを sycl::range と sycl::id で指定します。 「処理単位」と記載しましたが、 SYCL においてはこのことを work-item と呼びます。 なので、sycl::range で work-item の数を、sycl::id で各 work-item の ID を指定し、parallel_for で実際に各 work-item に実行させる処理内容を記載するということとなります。

ここで、 work-item の数がデバイスの最大スレッド数よりも小さいときは1バッチで処理が完了しますが、 work-item の数がデバイスの最大スレッド数よりも大きいときは複数のバッチで処理が行われることとなります。

このバッチは work-group と呼ばれます。

sycl::range と sycl::id

parallel_for の第 1 引数には sycl::range を指定します。sycl::range は work-item の数を示し、上記のプログラムでは N , すなわち 10 個の work-item を生成します。

今回のような1次元配列の例では特に sycl::range を用いる理由はありませんが、2次元配列や3次元配列の場合には sycl::range を用いることで、配列の次元数を明示的に示すことができます。 parallel_for の第 2 引数には SYCL カーネル関数が指定されますが、この関数オブジェクトの引数には sycl::id が渡されます。sycl::id は各 work-item に関する情報が渡され、 ID を確認することができます。

上記のプログラムでは get_linear_id() により、各 work-item の ID を取得し、配列 data の対応するインデックスにアクセスしています。 以上のことを cuda で置き換えると、 sycl::range は cuda の dim3 に相当し、 sycl::id は cuda の threadIdx や blockIdx などに相当する、ということができます。

少し複雑なプログラム

SYCL のプログラムをもう少し複雑なものにします。次のプログラムは、行列の乗算を SYCL で実装したものです。

#include <sycl/sycl.hpp>
#include <iostream>

int main() {
    constexpr int M = 2;
    constexpr int N = 3;
    constexpr int P = 4;

    // キューの作成
    sycl::queue q;

    // 統合共有メモリの作成
    auto A = sycl::malloc_shared<float>(M * N, q);
    auto B = sycl::malloc_shared<float>(N * P, q);
    auto C = sycl::malloc_shared<float>(M * P, q);

    // ホストで統合共有メモリの内容の初期化
    for (int i = 0; i < M * N; i++) {
        A[i] = i;
    }
    for (int i = 0; i < N * P; i++) {
        B[i] = i;
    }

    q.submit([&](sycl::handler &cgh) {
        // 2次元配列の parallel_for
        cgh.parallel_for(sycl::range<2>(M, P), [=](sycl::id<2> idx) {
            int i = idx[0]; // x の取得
            int j = idx[1]; // y の取得
            float sum = 0.0f;
            for (int k = 0; k < N; k++) {
                // デバイス上で統合共有メモリの操作
                sum += A[i * N + k] * B[k * P + j];
            }
            C[i * P + j] = sum;
        });
    });

    q.wait();

    for (int i = 0; i < M * P; i++) {
        std::cout << "C[" << i << "] = " << C[i] << std::endl;
    }

    return 0;
}

出力結果は以下の通りとなります。

C[0] = 20
C[1] = 23
C[2] = 26
C[3] = 29
C[4] = 56
C[5] = 68
C[6] = 80
C[7] = 92

最初に提示した簡単なプログラムと比較して、以下の点で違いがみられます。

  • sycl::malloc_shared を使用して、共有メモリ上に配列を確保している
  • 2次元の sycl::range を使用して、行列の乗算を並列化している
  • 2次元の sycl::id を使用して、行列の要素をアクセスしている
  • 並列実行の完了を待つために、q.wait() を使用している

以降は、これらの内容について説明します。

sycl::malloc_* : 統合共有メモリ（USM, Unified Shared Memory）

統合共有メモリー (USM) は、バッファー・プログラミング・モデルによるポインターベースの代替手段を提供します。USM は以下を可能にします。 • 割り当てアクセスをバッファーの代わりにポインターとすることにより、既存のコードベースへの統合が容易となり、割り当てられたメモリーへのポインター操作が完全にサポートされます。 • パフォーマンスとプログラマーの利便性を適切に選択するため、割り当ての所有権とアクセスを細かく制御できます。 • SYCL デバイスとホスト間で割り当てを自動的に移行する単純なプログラミング・モデルを実現します。

SYCL* 2020 より利用できる、 sycl::buffer と sycl::acesor の代替となる機能で、cuda では cudaMallocManaged に相当します。 sycl::buffer の内容は sycl::acesorによりアクセサを通してデバイスからアクセスする必要がありますが、USM ではポインタを通してホスト・デバイス双方からアクセスすることができます。

USM には、sycl::malloc_host、sycl::malloc_device、sycl::malloc_shared があり、 sycl::malloc_shared を用いることでホスト,デバイスの両方からアクセスできる単一の統合アドレスを使用することができます。

sycl::range : 並列実行の範囲

sycl::range は、並列実行の範囲を指定するために使用することはすでに示した通りです。

この例では、2次元配列へのアクセスのために2次元の sycl::range<2> を使用します。 最大3次元までを指定することができ、3次元の場合は sycl::range<3> となります。 cuda の dim3 と同様の使い方のため、第1引数には x 方向の大きさ、第2引数には y 方向の大きさ、第3引数には z 方向の大きさを指定します。

sycl::id : 並列実行のID

sycl::id は、並列実行の ID を取得するために使用することもすでに示した通りです。

sycl::range 同様に、最大3次元までを指定することができ、3次元の場合は sycl::id<3> となります。 cuda の threadIdx と似た使い方ですが、 ID の取得方法には違いがあり、sycl::id<3> idx としたときに、x 方向の ID は idx[0]、y 方向の ID は idx[1]、z 方向の ID は idx[2] と取得することができます。

wait() での同期

sycl::queue には wait() という関数があり、同期を取ることができます。 cuda での、cudaDeviceSynchronize() と同様の機能と考えることができます。

上記のプログラムでは、sycl::malloc_shared でのメモリ確保後、submit で処理の内容を記載し、この処理の完了を待つために wait() を使用しています。 この wait() の後では同期されているため、 sycl::malloc_shared の内容も submit で送信された処理が反映された内容となっています。 逆に、wait() の前では同期されていないため、sycl::malloc_shared の内容は submit で送信された処理が反映されているかは保証されません。

プログラムのビルド・実行

「少し複雑なプログラム」の内容に少し手を加え、配列のサイズを大きくし、時間計測コードを加えて、実際に複数のアーキテクチャで実行します。今回の実行環境は以下の2つを用意しました。

OS CPU GPU Memory
ubuntu22.04 i9-14900K NVIDIA RTX 3060Ti(CUDA 12.6) DDR5 64GB
ubuntu22.04 i9-12900K Intel Arc A770(OpenCL 3.0) DDR5 32GB

コンパイル時にGPUを指定できるので、この部分を変更してコンパイルし、各 PC で GPU を用いて実行するようにします。

実行するコードは以下の内容です。

#include <sycl/sycl.hpp>
#include <chrono>
#include <cmath>
#include <iostream>
#include <vector>

int main() {
    const auto repeat_num = 100;
    std::vector<float> time(repeat_num);

    constexpr int M = 1 << 10;
    constexpr int N = 1 << 10;
    constexpr int P = 1 << 10;

    sycl::queue q(sycl::gpu_selector_v);

    auto A = sycl::malloc_shared<float>(M * N, q);
    auto B = sycl::malloc_shared<float>(N * P, q);
    auto C = sycl::malloc_shared<float>(M * P, q);

    for (int i = 0; i < M * N; i++) {
        A[i] = i;
    }
    for (int i = 0; i < N * P; i++) {
        B[i] = i;
    }

    for (int n = 0; n < repeat_num; n++) {
        auto start = std::chrono::system_clock::now();

        q.submit([&](sycl::handler &cgh) {
            cgh.parallel_for(sycl::range<2>(M, P), [=](sycl::id<2> idx) {
                int i = idx[0];
                int j = idx[1];
                float sum = 0.0f;
                for (int k = 0; k < N; k++) {
                    sum += A[i * N + k] * B[k * P + j];
                }
                C[i * P + j] = sum;
            });
        });

        q.wait();

        auto end = std::chrono::system_clock::now();
        time[n] = std::chrono::duration_cast<std::chrono::microseconds>(end - start).count();
    }

    float sum = 0.0f;
    float mins = time[0];
    float maxs = time[0];
    for (int i = 0; i < repeat_num; i++) {
        sum += time[i];
        if (time[i] < mins) {
            mins = time[i];
        }
        if (time[i] > maxs) {
            maxs = time[i];
        }
    }
    float avg = sum / repeat_num;
    float diff_sum = 0.0f;
    for (int i = 0; i < repeat_num; i++) {
       diff_sum += pow(time[i] - avg, 2);
    }
    float sd = sqrt(diff_sum / repeat_num);
    
    std::cout << "Device: " << q.get_device().get_info<sycl::info::device::name>() << std::endl;
    std::cout << "Average time: " << avg << " usec" << std::endl;
    std::cout << "Min time: " << mins << " usec" << std::endl;
    std::cout << "Max time: " << maxs << " usec" << std::endl;
    std::cout << "Standard Deviation: " << sd << " usec" << std::endl;
    
    return 0;
}

SYCL 環境構築

今回は、 Intel oneAPI DPC++ を使用します。 環境構築方法は以前当ブログにて紹介いたしました、 SYCL 実装の性能を比較してみた および Build DPC++ toolchain with support for NVIDIA CUDA の内容を参照してください。

実行結果

利用する GPU を変更して実行した結果が以下の通りです。

今回の環境では、コンパイル時の -fsycl-targets オプションを用いて使用する GPU の変更を行いました。 RTX 3060 Ti を用いるようにした場合は、 -fsycl-targets=nvptx64-nvidia-cuda を追加し、 Arc A770 を用いるようにした場合は、このオプションを追加せずにコンパイルを行いました。

RTX 3060 Ti での実行時の出力は以下の通りです。

Device: NVIDIA GeForce RTX 3060 Ti
Average time: 2880.14 usec
Min time: 2824 usec
Max time: 5833 usec
Standard Deviation: 298.558 usec

Arc A770 での実行時の出力は以下の通りです。

Device: Intel(R) Arc(TM) A770 Graphics
Average time: 8173.93 usec
Min time: 8095 usec
Max time: 12936 usec
Standard Deviation: 480.527 usec

表にすると以下の通りとなります。

GPU 平均時間(μs) 最小時間(μs) 最大時間(μs) 標準偏差(μs)
GeForce RTX 3060 Ti 2880 2824 5833 299
Intel Arc A770 8174 8095 12936 481

RTX 3060 Ti , Arc A770 それぞれについて出力を確認でき、同一のコードから確かに各デバイスで実行ができていることが確認できます。

処理時間を確認しても、各 GPU の性能が反映されたものになっていることも確認できます。

まとめ

今回の記事では以下を行いました。

  • SYCL の簡単な構文・用語の説明
  • SYCL を使用したプログラムの説明
  • SYCL を使用したプログラムのビルド
    • 複数デバイスでの実行、処理時間計測
    • 同一のプログラムで複数デバイス上で実行できることの説明

参考文献

Tags

C++GPUSYCL

About Author

shinnosuke.takemoto

Leave a Comment

メールアドレスが公開されることはありません。 ※ が付いている欄は必須項目です

This site uses Akismet to reduce spam. Learn how your comment data is processed.

Recent Comments

Social Media