色々なデバイスでOpenCL してみた！

Transcript

1. 1 《発表タイトル》 色々なデバイスで OpenCL してみた！ 2025/5/30 (金) 本橋 弘臣 インターフェース 2025年6月号 オフ会
2. 2 《本職》 ㈱ＶＲＡＩＮ Ｓｏｌｕｔｉｏｎ(中央区晴海)に て、組み込みソフトウェア開発業務に 従事 《得意技術》 ・GPGPU (NVIDIA CUDA、AMD OpenCL) ・SIMD (Intel SSE/AVX、ARM Neon) ・Linux システムプログラミング ・RDMA (RoCE / iWARP) 高速通信技術 自己紹介
3. 3 デバイス C 言語 Open CL CUDA GPU RasPi5 (VC7) － ◦ × NVIDIA RTX-3070 － ◦ ◦ AMD GPU － ◦ ×(*1) Intel UHD 630 (*2) － ◦ × CPU RasPi5 (Cortex-A76) ◦ － － Intel Core i7-8700 ◦ ◦ － 今回取り上げるデバイス *1 AMD は CUDA 互換の HIP を提供している。が、not 100% compatible. *2 Intel Core i7-8700 内蔵 GPU
4. 4 《GPU ベンダーロックを回避できる》 ・同じソースコードが NVIDIA/AMD 両社の GPU で動作する。 ⇒GPU の相見積もりが可能になる。 ⇒コストダウンにつながるかも？？ 《CPU 処理でも性能最適化が可能》 ・自動的にマルチスレッド対応。 ・C コードを書くだけで、SIMD 最適化 コードが手に入る (かも)。 OpenCL って、おいしいの？
5. 5 ①各デバイスで OpenCL を利用する ための環境構築方法。 ②各デバイスでのベンチマークテスト 結果 ・アプリケーション題材： 音声データに 対する FIR フィルタ(*1) 処理 ・OS： Linux 本日ご紹介する内容 *1 FIR = Finite Impulse Response (有限インパルス応答) フィルタ ⇔ IIR = Infinite Impulse Response (無限インパルス応答)
6. 6 (おそらく皆さんご存知の)平滑化フィルタ 15220 10348 2108 -30000 -20000 -10000 0 10000 20000 30000 1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 入力データ 9225 -30000 -10000 10000 30000 1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 平滑化データ 平均値算出 (15220+10348+2108)÷3 = 9225.33 ⇒注目データを中心に前後 3 点のデータの平 均値を得る事で、LPF 的な効果が得られている。 16bit 音声データ ・実は平滑化フィ ルタも FIR フィ ルタの一つ。 ・フィルタ Tap 数 = 3 ・フィルタ係数： { 1/3, 1/3, 1/3 } 注目データ
7. 7 ◦フィルタ係数のみを変化させるだけ で、LPF/HPF/BPF など好きな周波 数特性を実現できる！ ◦波形の崩れが少ない素直な特性の フィルタが得られる。 ×スパッと切れる急峻なカットオフ特 性を持つフィルタを作ろうとすると、 フィルタの Tap 数が多くなりがち。 ⇒演算量が多いので GPU 向き。 FIR フィルタの特徴
8. 8 FIR フィルタの設計方法 ・「FIR フィルタの設計」＝「フィルタ係 数の算出 」 (Tap 数の決定を含む) ・具体的にどうやってフィルタ係数を算 出するか？ ⇒教科書を読んで頑張る。 ⇒フィルタ設計ツールを利用する。 ⇒適当な Web サイトを利用する。 ・http://dsp.jpn.org/dfdesign ・https://hs-soft.exout.net/
9. 9 BPF フィルタの設計結果 -100 -90 -80 -70 -60 -50 -40 -30 -20 -10 0 10 0 1000 2000 3000 4000 5000 6000 7000 8000 ゲイン [dB] 周波数 [Hz] BPF (500-3500Hz、FIR 97Tap、Hanning 窓) -100 -90 -80 -70 -60 -50 -40 -30 -20 -10 0 10 0 1000 2000 3000 4000 5000 6000 7000 8000 ゲイン [dB] 周波数 [Hz] BPF (500-3500Hz、FIR 199Tap、Hanning 窓) a01 = 0 a02 = -5.706805009773E-6 a03 = -2.0039056830324E-5 ： a96 = -5.706805009773E-6 a97 = 0 得られたフィルタ係数 -0.05 0 0.05 0.1 0.15 1 11 21 31 41 51 61 71 81 91 フィルタ係数 ↓ グラフ化すると ⇒フィルタの Tap 数が長い ほど、カットオフ特性が良 好なフィルタになる。
10. 10 BPF フィルタの設計結果 -100 -90 -80 -70 -60 -50 -40 -30 -20 -10 0 10 0 1000 2000 3000 4000 5000 6000 7000 8000 ゲイン [dB] 周波数 [Hz] BPF (500-3500Hz、FIR 97Tap、Hanning 窓) a01 = 0 a02 = -5.706805009773E-6 a03 = -2.0039056830324E-5 ： a96 = -5.706805009773E-6 a97 = 0 各種デバイスでのベンチ マークテストは、この Tap=97 の FIR フィルタ演 算処理で実施した。 -0.05 0 0.05 0.1 0.15 1 11 21 31 41 51 61 71 81 91 フィルタ係数 ↓ グラフ化すると 得られたフィルタ係数
11. 11 最適化版 GPU コードの開発方法 作業内容 ① 普通の C コードを作成する。処理効率よりも 間違いなく正しい演算を行うことを重視する。 ② ①のコードを元に、C 言語レベルで処理性 能を最適化する。 ③ ②のコードを GPU に移植する。 ④ ③のコードを GPU 向けに最適化する。 ※ ①のプログラムの出力を "正解データ"(あるいは期 待値) として保存しておき、②～④の出力が正解と一致 することを常に確認しながら作業を進めるとよい。
12. 12 入出力データ形式 《入力音声データ》 LEFT RIGHT LEFT RIGHT LEFT RIGHT LEFT RIGHT 4Byte(float) *1 音声１サンプル 《フィルタ係数》 例： Tap数=5 0.1 0.2 《0.4》 0.2 0.1 《出力音声データ》 LEFT RIGHT LEFT RIGHT LEFT RIGHT LEFT RIGHT 4Byte(float) *1 FP32 の音声データは -1.0 < (サンプル値) < 1.0 の範囲で正規化され ている。ちなみに 16bit (int) では -32768～32767 の範囲となる。
13. 13 ①普通の C コードを作成する 《入力音声データ》 LEFT RIGHT LEFT RIGHT LEFT RIGHT LEFT RIGHT 4Byte(float) // 音声１サンプル分のデータ型 typedef struct sample_t { float left; float right; } SAMPLE; SAMPLE *inbuf; int samples = nnnnn; // 音声データのサンプル数 inbuf = malloc(sizeof(SAMPLE) * samples); // 入力音声データを inbuf に読み込む。 :
14. 14 ①普通の C コードを作成する 《フィルタ係数》 0.1 0.2 0.4 0.2 0.1 4Byte(float) float filter[5] = { 0.1f, 0.2f, 0.4f, 0.2f, 0.1f };
15. 15 ①普通の C コードを作成する 《FIR フィルタ処理》 // FIR フィルタの出力データ１個を計算する関数 SAMPLE calc_inner_product(const SAMPLE *inbuf, float *filter, int tap) { int i; SAMPLE sum = { 0.0f, 0.0f }; for (i = 0; i < tap; i++) { sum.left += inbuf[i].left * filter[i]; sum.right += inbuf[i].right * filter[i]; } return sum; } 実はこのコードでは、あまり性能が出ません。 それはなぜなのか？？ ⇒こういう時はメモリアクセスに注目して図を書くとよい。
16. 16 なぜ①のコードだと性能が出ないか 入力データ 0 0 0 0 L0 R0 L1 R1 L2 R2 注目データ フィルタ係数 0.1 0.2 0.4 0.2 0.1 乗算結果 (L-CH) 0 0 0.4・L0 0.2・L1 0.1・L2 × × × × × 入力データ 0 0 0 0 L0 R0 L1 R1 L2 R2 フィルタ係数 0.1 0.2 0.4 0.2 0.1 乗算結果 (R-CH) 0 0 0.4・R0 0.2・R1 0.1・R2 × × × × × L-CH と R-CH を別々に計算する羽目になるのは何故？ ⇒フィルタ係数がモノラルで、入力データと合わないから L-CH R-CH L-CH
17. 0 0 0 0 0.4・L0 0.4・R0 0.2・L1 0.2・R1 0.1・L2 0.1・R2 0.1 0.1 0.2 0.2 0.4 0.4 0.2 0.2 0.1 0.1 17 フィルタ係数をステレオ化する 入力データ 0 0 0 0 L0 R0 L1 R1 L2 R2 注目データ (音声データ 1 サンプル分) フィルタ係数 乗算結果 L-CH R-CH L-CH × × × × × × × × × × 4B×4 = 16Byte = 128bit ・Intel SSE 命令を使用すれば、一度に計算できる！！ ・ARM Neon 命令を使用すれば、 〃 。 4B×8 = 32Byte = 256bit ・Intel AVX 命令を使用すれば、一度に計算可能。
18. 18 最適化版 GPU コードの開発方法 作業内容 ① 普通の C コードを作成する。処理効率よりも 間違いなく正しい演算を行うことを重視する。 ② ①のコードを元に、C 言語レベルで処理性 能を最適化する。 ③ ②のコードを GPU に移植する。 ④ ③のコードを GPU 向けに最適化する。 ※ ①のプログラムの出力を "正解データ"(あるいは期 待値) として保存しておき、②～④の出力が正解と一致 することを常に確認しながら作業を進めるとよい。
19. 19 ② C 言語レベルで最適化 // FIR フィルタの出力データ１個を計算する関数 SAMPLE calc_inner_product(const SAMPLE *inbuf, float *filter, int tap) { for (i = 0; i < tap; i++) { sum.left += inbuf[i].left * filter[i]; sum.right += inbuf[i].right * filter[i]; } // FIR フィルタの出力データ１個を計算する関数 (fp32x2) SAMPLE calc_inner_product(const SAMPLE *inbuf, SAMPLE *filter, int tap) { for (i = 0; i < tap; i++) { sum.left += inbuf[i].left * filter[i].left; sum.right += inbuf[i].right * filter[i].right; } 最適化 さらに for ループ 1 回当たりの計算量を 2 倍に増やし たプログラムも作成してみました。(fp32x4_…)
20. 20 処理時間測定結果：プログラム② C 言語レベルの最適化はかなり効果があった。 ⇒この最適化は OpenCL Kernel コードと非常に合う。 ・測定条件： 入力データの長さは 2分32秒 (fs=44100)
21. 21 最適化版 GPU コードの開発方法 作業内容 ① 普通の C コードを作成する。処理効率よりも 間違いなく正しい演算を行うことを重視する。 ② ①のコードを元に、C 言語レベルで処理性 能を最適化する。 ③ ②のコードを GPU に移植する。 ④ ③のコードを GPU 向けに最適化する。 ※ ①のプログラムの出力を "正解データ"(あるいは期 待値) として保存しておき、②～④の出力が正解と一致 することを常に確認しながら作業を進めるとよい。
22. // FIR フィルタの出力データ１個を計算する関数 (fp32x2) SAMPLE calc_inner_product(const SAMPLE *inbuf, SAMPLE *filter, int tap) { for (i = 0; i < tap; i++) { sum.left += inbuf[i].left * filter[i].left; sum.right += inbuf[i].right * filter[i].right; } 22 ③ C コードを GPU に移植する OpenCL Kernel コード化 __kernel void kernel_func(__global float2 *inbuf, __global float2 *outbuf, __global float2 *filter) { float2 sum2; uint x_gid = get_global_id(0); inbuf += x_gid; sum2 = inbuf[ 0] * filter[ 0] + : inbuf[ 96] * filter[ 96]; outbuf[x_gid] = sum2; } ループを展開 (アンロール)
23. 23 最適化版 GPU コードの開発方法 作業内容 ① 普通の C コードを作成する。処理効率よりも 間違いなく正しい演算を行うことを重視する。 ② ①のコードを元に、C 言語レベルで処理性 能を最適化する。 ③ ②のコードを GPU に移植する。 ④ ③のコードを GPU 向けに最適化する。 ※ ①のプログラムの出力を "正解データ"(あるいは期 待値) として保存しておき、②～④の出力が正解と一致 することを常に確認しながら作業を進めるとよい。
24. 24 VideoCore VII (RasPi-5) の内部構造 GPU (Broadcom VideoCore VII) Slice Slice Slice Slice Slice L1 Cache (I-Cache) SFU QPU QPU QPU QPU L2 Cache (命令/データ) LPDDR4 Memory QPU (CU) PC (Program Counter) PE PE PE PE Processor Element 加算系 命令用 演算 パイプ 乗算系 命令用 演算 パイプ ・加算(int/float) ・bit 演算 ・8bit SIMD ・乗算 (int/float) ・8bit SIMD 各パイプは *一度に* (32bit×4) のデータに 対する演算 が可能 拡大
25. 25 ④ GPU コードを最適化する ・今回は時間の関係で、これ以上の 最適化は実施しなかった。 ・ただ次の最適化の方向性は、カーネ ルコード内での演算単位を float2 から float4 に変更するだと考える。
26. 26 今回作成したプログラムの仕様 項目 仕様 入力データ RAW、float(FP32)、STEREO (L→R) 出力データ 〃 データ入力元 標準入力 データ出力先 標準出力 又は ファイル 書式 (コマン ドライン引数) $ ./main <フィルタ係数ファイル> <CPU 実行結果出力先ファイル> [<GPU 実行結果出力先ファイル>] < <入力音声ファイル>
27. 27 プログラム実行例 (RasPi-5) $ sox "input.wav" -t raw -e float -b 32 input-fp32.raw // input.wav(16bit/int) を input-fp32.raw(32bit/float) に変換 $ ./main bpf97-500-3500.txt out-cpu-fp32.raw out-gpu-fp32.raw < input-fp32.raw filter-tap=97 samples=6731136 [CPU] bench_fir_filter_cpu() Time=449.4[ms] BW=23.4[GB/s] [CPU] bench_fir_filter_cpu() Time=400.1[ms] BW=26.2[GB/s] [CPU] bench_fir_filter_cpu() Time=400.1[ms] BW=26.2[GB/s] out-cpu-fp32.raw saved. [Host] device_name=V3D 7.1.7.0 // Broadcom VideoCoreVII (RP5) [Host] clEnqueueNDRangeKernel() gws=(6731136) lws=(256) [GPU] bench_fir_filter_opencl() Time=673.27[ms] BW=15.6[GB/s] [GPU] bench_fir_filter_opencl() Time=673.21[ms] BW=15.6[GB/s] [GPU] bench_fir_filter_opencl() Time=712.86[ms] BW=14.7[GB/s] out-gpu-fp32.raw saved. $ sox --no-dither -t raw -e float --bits 32 --rate 44100 -- channels 2 out-gpu-fp32.raw -e signed-integer --bits 16 out- gpu-int16.wav // out-gpu-fp32.raw を out-gpu-int16.wav に変換
28. 28 ソースコード構成 ファイル名 内容 common.h プログラム全体の共通ヘッダ cpu-bench.c CPU で FIR フィルタ処理のベンチマークテス トを実行する関数 cpu-bench.h cpu-bench.c の公開関数定義 kernel.cl OpenCL カーネルコード main.c メイン関数 mk.sh ビルド用シェルスクリプト opencl-bench.c GPU で FIR フィルタ処理のベンチマークテスト を実行する関数 opencl-bench.h gpu-bench.c の公開関数定義
29. 29 プログラム動作設定パラメータ ファイル名 コード 意味 common.h L4 #define USE_STEREO_FILTER ステレオ形式のフィル タ係数データを使用 するかどうか (CPU 演算時向け) opencl-bench.c L223 set_work_size(lws, 256, 0); ローカルワークサイズ kernel.cl L1 #define KERNEL_VEC_SIZE SIMD ベクタ長 (但し、 KERNEL_VEC_SIZE=4 はバグがあり、正常 に動作しない)
30. 30 ベンチマークテスト結果：RasPi-5(GPU) ⇒CPU 処理時の性能がかなり速い。 ⇒速さの秘訣は、gcc の自動ベクトル化が非常にうまく 働いていたから。 ⇒その結果、ARM-Neon 命令を利用して FP32x4 を 一度に計算する、(おそらく ARM では) 最速のアセンブ ラコードが生成された。
31. 31 《前提条件》 ・OS： Ubuntu 24.04 (x86_64) NVIDIA GPU で OpenCL を利用する方法 # wget https://developer.download.nvidia.com/compute/cuda/ repos/ubuntu2404/x86_64/cuda-keyring_1.1-1_all.deb # dpkg -i cuda-keyring_1.1-1_all.deb # apt update # apt install cuda-drivers cuda-toolkit-12-6 // 少し枯れたver $ vi ~/.bashrc // ファイル末尾に以下の３行を追加 -- # CUDA Toolkit export PATH=$PATH:/usr/local/cuda/bin export LD_LIBRARY_PATH=$LD_LIBRARY_PATH:/usr/local/cuda/lib64 -- $ nvidia-smi $ clinfo (*1) *1 clinfo コマンド実行時に NVIDIA GPU が表示されれば OK。
32. 32 《前提条件》 ・OS： Ubuntu 24.04 (x86_64) NVIDIA GPU で OpenCL を利用する方法 # wget https://developer.download.nvidia.com/compute/cuda/ repos/ubuntu2404/x86_64/cuda-keyring_1.1-1_all.deb # dpkg -i cuda-keyring_1.1-1_all.deb # apt update # apt install cuda-drivers cuda-toolkit-12-6 // 少し枯れたver $ vi ~/.bashrc // ファイル末尾に以下の３行を追加 -- # CUDA Toolkit export PATH=$PATH:/usr/local/cuda/bin export LD_LIBRARY_PATH=$LD_LIBRARY_PATH:/usr/local/cuda/lib64 -- $ nvidia-smi $ clinfo (*1) *1 clinfo コマンド実行時に NVIDIA GPU が表示されれば OK。
33. 33 ベンチマークテスト結果：RTX-3070 ⇒RasPi-5 の 500 倍位速い。さすがは本物の GPU!! ⇒OpenCL と CUDA の性能差は大きくない。 ⇒独自にカーネルコードを書くなら、OpenCL で書けば よいのでは？(GPU ベンダーフリーになるし)
34. 34 ・以下のリンク先のページで OpenCL 対応の GPU ドライバを入手することができる。 ・Radeon Software for Linux 25.10.1 →https://www.amd.com/en/resources/supp ort-articles/release-notes/RN-AMDGPU- UNIFIED-LINUX-25-10-1.html AMD GPU で OpenCL を利用する方法
35. 35 AMD：サポート対象 GPU リスト
36. 36 ・最新の Intel Graphics OpenCL ドライバはここで手 に入る。 →https://github.com/intel/compute-runtime ・Broadwell(Core i-5xxx)～Ice Lake(10xxx番台) は Legacy システムと位置付けられている。 ・Legacy システムがサポートされるのは、Intel Graphics Compute Runtime の v24.35 まで。 →https://github.com/intel/compute- runtime/releases/tag/24.35.30872.22 Intel GPU で OpenCL を利用する方法
37. 37 Intel Graphics Compute Runtime v24.35 のインストール # cd # mkdir -p tmp/neo # cd tmp/neo # wget https://github.com/intel/intel-graphics-compiler/releases/download/igc- 1.0.17537.20/intel-igc-core_1.0.17537.20_amd64.deb # wget https://github.com/intel/intel-graphics-compiler/releases/download/igc- 1.0.17537.20/intel-igc-opencl_1.0.17537.20_amd64.deb # wget https://github.com/intel/compute-runtime/releases/download/24.35.30872.22/intel- level-zero-gpu-dbgsym_1.3.30872.22_amd64.ddeb # wget https://github.com/intel/compute-runtime/releases/download/24.35.30872.22/intel- level-zero-gpu-legacy1-dbgsym_1.3.30872.22_amd64.ddeb # wget https://github.com/intel/compute-runtime/releases/download/24.35.30872.22/intel- level-zero-gpu-legacy1_1.3.30872.22_amd64.deb # wget https://github.com/intel/compute-runtime/releases/download/24.35.30872.22/intel- level-zero-gpu_1.3.30872.22_amd64.deb # wget https://github.com/intel/compute-runtime/releases/download/24.35.30872.22/intel- opencl-icd-dbgsym_24.35.30872.22_amd64.ddeb # wget https://github.com/intel/compute-runtime/releases/download/24.35.30872.22/intel- opencl-icd-legacy1-dbgsym_24.35.30872.22_amd64.ddeb # wget https://github.com/intel/compute-runtime/releases/download/24.35.30872.22/intel- opencl-icd-legacy1_24.35.30872.22_amd64.deb # wget https://github.com/intel/compute-runtime/releases/download/24.35.30872.22/intel- opencl-icd_24.35.30872.22_amd64.deb # wget https://github.com/intel/compute- runtime/releases/download/24.35.30872.22/libigdgmm12_22.5.0_amd64.deb # dpkg -i *.deb $ clinfo
38. 38 ベンチマークテスト結果： Intel UHD 630 ⇒対 RasPi-5 では 14 倍位速いが、RTX-3070 と比べ ると 1/40 位の性能。 ⇒CPU+普通の C プログラム+シングルスレッド実行と 比べると 5 倍位速い。 ⇒つまり本物の GPU とは比べ物にならない位遅いが、 CPU と比べるとそこそこ速い。
39. 39 ・Intel CPU 向けの OpenCL ランタイムに関する情報 はここで手に入る。 ・Intel CPU Runtime for OpenCL Applications with SYCL support →https://www.intel.com/content/www/us/en/de veloper/articles/technical/intel-cpu-runtime-for- opencl-applications-with-sycl-support.html ・Intel oneAPI Toolkits Installation Guide for Linux OS →https://www.intel.com/content/www/us/en/do cs/oneapi/installation-guide-linux/2023- 0/apt.html#GUID-560A487B-1B5B-4406-BB93- 22BC7B526BCD Intel CPU で OpenCL を利用する方法
40. 40 OpenCL runtime for Intel CPU をインストール # wget -O- https://apt.repos.intel.com/intel-gpg-keys/GPG-PUB-KEY- INTEL-SW-PRODUCTS.PUB | gpg --dearmor > /usr/share/keyrings/oneapi- archive-keyring.gpg # vi /etc/apt/sources.list.d/oneAPI.list // "--" ～ "--" で囲まれた ２行を入力 -- deb [signed-by=/usr/share/keyrings/oneapi-archive-keyring.gpg] https://apt.repos.intel.com/oneapi all main $ clinfo -- # apt update # apt install intel-oneapi-runtime-opencl $ clinfo ⇒同一のシステムに複数の OpenCL ランタイムをイン ストールすると、clinfo コマンド実行時に複数の OpenCL デバイスが表示されるようになる。
41. 41 ベンチマークテスト結果： Core i7-8700 ⇒Intel 内蔵 Graphic と比較すると 90% 位の性能。 ⇒CPU+普通の C プログラム+シングルスレッド実行と 比べると 4 倍位速い。 ⇒CPU でお手軽に実行速度を最適化する手段として、 OpenCL が利用できるということ。
42. 42 ・/etc/OpenCL/vendors ディレクトリ以下のファイルを リネームすることで、使用する OpenCL デバイスを限 定することができる。 使用する OpenCL デバイスを選択する方法 # cd /etc/OpenCL/vendors # ls -1 intel64.icd // Intel CPU 用 intel_legacy1.icd.x // Intel GPU 用 nvidia.icd.x // NVIDIA GPU 用 ⇒*.icd ファイルを適当な別の拡張子にリネームしてお くと、そのデバイスが使用不可状態となる。
43. 43 OpenCL ベンチマークテスト結果まとめ ⇒本物の GPU の処理性能は桁違いに物凄い。 ⇒Intel CPU と Intel GPU の性能が近いのは、メモリ バンド幅がボトルネックになっているのかもしれない。 ⇒プロセッサコア当たりの性能は Intel CPU がダントツ。 ⇒RasPi の GPU は世代更新と共に少しずつ速くなって きているとはいえ、PC 向け GPU と比べると遅い。
44. 44 ①OpenCL でプログラムを書くと、そのコードは 様々なデバイス上でそのまま動作するという 話は本当だった。 ②NVIDIA GPU 向けに独自にカーネルコード を開発する必要が生じた場合には、OpenCL を選択してもよいのかもしれない。 ③昔の常識では NVIDIA GPU では性能的に CUDA の一択だった。しかし、OpenCL 利用 時のパフォーマンス問題はかなり改善されて いる模様。 その他の考察