CUDA高速化セミナーvol.1 ～画像処理アルゴリズムの高速化～

2022年5月27日に開催された「CUDA高速化セミナーvol.1　～画像処理アルゴリズムの高速化～」の当日資料です。

Transcript

1. Fixstars Group www.fixstars.com Copyright © Fixstars Group Copyright © Fixstars Group CUDA 高速化セミナー vol.1 画像処理アルゴリズムの高速化
2. Fixstars Group www.fixstars.com Copyright © Fixstars Group CUDA高速化セミナーをシリーズ化 2 「いまさら聞けないCUDA高速化」が好評につきシリーズ化 CUDA高速化セミナー • Vol.1 画像処理アルゴリズムの高速化（いまさら聞けないCUDA高速化の実践特化編） • Vol.2 CUDAアーキテクチャの進化
3. Fixstars Group www.fixstars.com Copyright © Fixstars Group 発表者紹介 3 • 冨田 明彦（とみた あきひこ） ソリューションカンパニー 営業企画執行役 2008年に入社。金融、医療業界において、 ソフトウェア高速化業務に携わる。その 後、新規事業企画、半導体業界の事業を 担当し、現職。 • 上野 晃司（うえの こうじ） ソリューション第一事業部 エグゼクティブエンジニア 2016年に入社。学生時代から続けている スパコンのベンチマークGraph500の 「京」「富岳」向け最適化の他、CUDA やOpenCLを使った画像処理高速化を担 当。
4. Fixstars Group www.fixstars.com Copyright © Fixstars Group 本日のAgenda フィックスターズの紹介 (15分) • 会社紹介 • 本ウェビナーに該当する、高速化サービスにおける開発プロセス CUDA高速化の復習 (10分) CUDA高速化の実践：ガウシアンフィルタ(30分) • CUDA化 • データ転送 • 実装方法による性能の比較 • RGB画像への対応 Q&A / 告知 4
5. Fixstars Group www.fixstars.com Copyright © Fixstars Group Copyright © Fixstars Group フィックスターズのご紹介
6. Fixstars Group www.fixstars.com Copyright © Fixstars Group フィックスターズの強み コンピュータの性能を最大限に引き出す、ソフトウェア高速化のエキスパート集団 低レイヤ ソフトウェア技術 アルゴリズム 実装力 各産業・研究 分野の知見 6
7. Fixstars Group www.fixstars.com Copyright © Fixstars Group 性能に関する課題 7 生産効率の向上 • より短時間で欠陥検出 • より安価なハードで 安全性の向上 • より精度の高い物体検出 • より低消費電力なハードで
8. Fixstars Group www.fixstars.com Copyright © Fixstars Group ソフトウェア高速化サービス (概要) お客様のソースコードをご提供いただき、 最適化やアルゴリズムの改良を施して高速化してお返しします 当社 お客様 オリジナルソースコードのご提供 高速化したソースコード コンサルティング 高速化 サポート 要件分析 先行研究等の調査 アルゴリズムの改良・開発 ハードウェアへの最適化 実製品への組込み支援 レポートやコードへのQ&A 8
9. Fixstars Group www.fixstars.com Copyright © Fixstars Group ソフトウェア高速化サービス 様々な領域でソフトウェア高速化サービスを提供しています 大量データの高速処理は、お客様の製品競争力の源泉となっています ・NAND型フラッシュメモリ向けファー ムウェア開発 ・次世代AIチップ向け開発環境基盤開発 Semiconductor ・デリバティブシステムの高速化 ・HFT(アルゴリズムトレード)の高速化 Finance ・自動運転の高性能化、実用化 ・次世代パーソナルモビリティの研究開発 Mobility ・ゲノム解析の高速化 ・医用画像処理の高速化 ・AI画像診断システムの研究開発 Life Science ・Smart Factory化支援 ・マシンビジョンシステムの高速化 Industrial 9
10. Fixstars Group www.fixstars.com Copyright © Fixstars Group 画像処理・アルゴリズム開発サービス 出展：https://www.cs.toronto.edu/~frossard/post/vgg16/ • お客様の課題 • 高度な画像処理や深層学習等のアルゴリズム開発を行える人材が社内に限られている • 考案中のアルゴリズムで機能要件は満たせそうだが、ターゲット機器上で性能要件まで クリアできるか不安 • 製品化に結びつくような研究ができていない • 弊社の支援内容 • 課題に応じたアルゴリズム調査 • 深層学習ネットワーク精度改善、推論高速化手法調査 • 論文調査、実装 10
11. Fixstars Group www.fixstars.com Copyright © Fixstars Group AI・深層学習関連サービス • ディープラーニングの包括的開発技術 • ネットワーク設計からターゲットデバイスでの高速化のノウハウ • 大規模システムからエッジコンピューティングまでの開発実績 ネットワーク設計 データの前処理、データ拡張 精度改善 分散処理による学習高速化 各種DLフレームワーク クラウド・サーバ エッジ モデル圧縮 - 量子化 - 枝刈り - 蒸留 ターゲットデバイスへの ポーティング及び推論高速化 ▪ ARM, GPU, DSP ▪ SIMD,NEON,CUDA,TensorRT 11
12. Fixstars Group www.fixstars.com Copyright © Fixstars Group GPU向け高速化サービス • お客様の課題 • GPU 高速化の知見がない • 自力で GPU に乗せてみたものの望む性能が出ない • 弊社の支援内容 • GPU 高速化に関するコンサルティング • ボトルネック調査、GPU プログラムの高速化 • CPU/GPU が混在するヘテロジニアス環境での最適化 10～150 倍の 高速化事例あり 12
13. Fixstars Group www.fixstars.com Copyright © Fixstars Group 本ウェビナーの対象プロセス 13 画像処理アルゴリズムを題材に 高速化の実践例をご紹介 要件分析 研究調査・アルゴリズム実装 高速化 / 評価 アルゴリズム改善 / 評価 品質確保
14. Fixstars Group www.fixstars.com Copyright © Fixstars Group • 社内大学 • プログラミングコンテスト • 勉強会 • 各種コンテストへの参加 • 勉強会 • 論文・学会発表 • 社内向け • 社外向け よりよいサービスのご提供を目指して • 組込み開発 • アルゴリズム開発 • AI・深層学習 • 組合せ最適化 各種高速化サービス 技術力強化 14
15. Fixstars Group www.fixstars.com Copyright © Fixstars Group Copyright © Fixstars Group CUDA高速化入門復習
16. Fixstars Group www.fixstars.com Copyright © Fixstars Group なぜGPUを使うのか • CPUと比べて • ピーク性能の高さ • 電力効率の良さ 16 浮動小数点数演算性能 メモリバンド幅 TDP 価格 CPU: AMD Ryzen 9 5950X 2.25※ [TFLOPS] 51.2 [GB/s] 105 [W] ¥70,000~ GPU: NVIDIA GeForce RTX 3070 20.31 [TFLOPS] 448.0 [GB/s] 220 [W] ¥83,000~ • その他のアクセラレータと比べて • 入手性・価格性能比の良さ • プログラミングの容易さ ※CPUは全コア4.4GHzで動作したと仮定
17. Fixstars Group www.fixstars.com Copyright © Fixstars Group なぜGPUが速いのか • 並列計算に特化した構成 • 大量のコア・演算器 • CPU: AMD EPYC 7763: 64 Cores, 32 FLOPs/Core/cycle • GPU: NVIDIA A100: 108 SMs, 128 FLOPs/SM/cycle • バス幅の広い広帯域メモリ • もちろん弱点もある • 並列に処理できない問題には弱い • 最大メモリ容量が小さい 17
18. Fixstars Group www.fixstars.com Copyright © Fixstars Group ホストメモリとデバイスメモリ • CPUとGPUはそれぞれがメモリを持っている • 目的に応じて適切なほうを利用する • 必要に応じて片方から他方へデータをコピーする 18 CPU GPU ホストメモリ (DDR) ～200 GB/s デバイスメモリ (GDDR/HBM) ～2000 GB/s ～20 GB/s
19. Fixstars Group www.fixstars.com Copyright © Fixstars Group スレッドの階層構造 • CUDAではスレッド間に階層構造がある • 近いスレッド同士はより密に通信・同期を行うことができる 19 Warp (32T) … Thread Block (～1024T) … Grid …
20. Fixstars Group www.fixstars.com Copyright © Fixstars Group • メモリにも階層構造がある • おおむねスレッドの階層構造と対応 メモリの階層構造 20 Global Memory Grid Thread Block Thread Registers Local Memory Shared Memory Constant Memory
21. Fixstars Group www.fixstars.com Copyright © Fixstars Group Copyright © Fixstars Group プロファイラ
22. Fixstars Group www.fixstars.com Copyright © Fixstars Group プロファイラー • プロファイラーは性能を分析するツール • CUDAのボトルネック解析や最適化に必須 • VoltaまでのGPUなら • NVIDIA Visual Profiler • Turing世代以降のGPUの場合 • NVIDIA Nsight Systems • NVIDIA Nsight Compute 22
23. Fixstars Group www.fixstars.com Copyright © Fixstars Group NVIDIA Nsight Systems • タイムラインの表示をサポート 23
24. Fixstars Group www.fixstars.com Copyright © Fixstars Group NVIDIA Nsight Compute • カーネルプロファイラをサポート 24
25. Fixstars Group www.fixstars.com Copyright © Fixstars Group Copyright © Fixstars Group CUDA高速化の実践 ガウシアンフィルタ CUDA化
26. Fixstars Group www.fixstars.com Copyright © Fixstars Group 本日説明するコード • ↓ここにあります • https://github.com/fixstars/CudaOptimizeSample/blob/master/CudaOptimize Sample/kernel.cu 26
27. Fixstars Group www.fixstars.com Copyright © Fixstars Group ガウシアンフィルタ ⊗ カーネル 27
28. Fixstars Group www.fixstars.com Copyright © Fixstars Group CPU版 void GaussianKernelCPU(const uint8_t *src, uint8_t *dst, int width, int height, int step) { const float filter[5][5] = { … }; for (int y = 0; y < height; ++y) { for (int x = 0; x < width; ++x) { float sum = 0; for (int dy = 0; dy < 5; ++dy) { for (int dx = 0; dx < 5; ++dx) { sum += filter[dy][dx] * src[(x + dx) + (y + dy) * step]; } } dst[x + y * step] = (int)(sum + 0.5f); } } } カーネル（値はソースコード参照） 画像xyループ カーネルxyループ とりあえず単純な問題から説明するため 1chの画像を対象とする 28
29. Fixstars Group www.fixstars.com Copyright © Fixstars Group ガウシアンフィルタCUDA化 スレッド割り当て ブロック (0,0) ブロック (0,1) ブロック (0,2) ブロック (0,3) ブロック (0,4) ブロック (1,0) ブロック (1,1) ブロック (1,2) ブロック (1,3) ブロック (1,4) ブロック (2,0) ブロック (2,1) ブロック (2,2) ブロック (2,3) ブロック (2,4) ブロック (3,0) ブロック (3,1) ブロック (3,2) ブロック (3,3) ブロック (3,4) 32 32 • 1スレッドが出力1ピク セルを担当 • ブロックの最大スレッ ド数は1024なので、１ ブロック 32x32(=1024スレッ ド)に設定 • 画像全体を覆うように ブロックを起動する 29
30. Fixstars Group www.fixstars.com Copyright © Fixstars Group ガウシアンフィルタCUDA化 単純移植カーネル __global__ void GaussianKernelSimple(const uint8_t *src, uint8_t *dst, int width, int height, int step) { const float filter[5][5] = { … }; int x = blockIdx.x * blockDim.x + threadIdx.x; int y = blockIdx.y * blockDim.y + threadIdx.y; if (x < width && y < height) { float sum = 0; for (int dy = 0; dy < 5; ++dy) { for (int dx = 0; dx < 5; ++dx) { sum += filter[dy][dx] * src[(x + dx) + (y + dy) * step]; } } dst[x + y * step] = (int)(sum + 0.5f); } } カーネル カーネルxyループ 画像xyループがなくなって、 代わりにスレッドIDになった 画像からはみ出すのを防ぐ 30
31. Fixstars Group www.fixstars.com Copyright © Fixstars Group ガウシアンフィルタCUDA化 カーネル呼び出し部分 cv::Mat GaussianFilterGPUSimple(cv::Mat src) { int width = src.cols, height = src.rows; uint8_t *dev_src, *dev_dst; ck(cudaMalloc((void**)&dev_src, width * height * sizeof(uint8_t))); ck(cudaMalloc((void**)&dev_dst, width * height * sizeof(uint8_t))); ck(cudaMemcpy(dev_src, src.data, width * height * sizeof(uint8_t), cudaMemcpyHostToDevice)); dim3 threadsPerBlock(32, 32); dim3 numBlocks((width + threadsPerBlock.x - 1) / threadsPerBlock.x, (height + threadsPerBlock.y - 1) / threadsPerBlock.y); GaussianKernelSimple <<<numBlocks, threadsPerBlock >>>(dev_src, dev_dst, width - 4, height - 4, width); cv::Mat dst(src.rows, src.cols, src.type()); ck(cudaMemcpy(dst.data, dev_dst, width * height * sizeof(uint8_t), cudaMemcpyDeviceToHost)); ck(cudaFree(dev_src)); ck(cudaFree(dev_dst)); return dst; } メモリ確保 カーネル起動 CPUに出力デー タを転送 入力データを GPUに転送 31
32. Fixstars Group www.fixstars.com Copyright © Fixstars Group ガウシアンフィルタCUDA化 単純移植カーネル • 25倍くらいになった • CPU(マルチスレッド)は OpenMPで単純に並列化した 実装 510 67 2.605 0 100 200 300 400 500 600 CPU シングルスレッド CPU マルチスレッド CUDA ガウシアンフィルタ計算時間 (ms) 計測環境 CPU: Core i7-8700 3.2GHz (6コア 12スレッド) GPU: GeForce RTX 2060 OS: Windows 10 計測条件 6720x4480の画像(グレースケール)を処理 計算時間のみで、データ転送やメモリ確保などの 時間を含めず 25.7倍 • ただし、データ転送も含めると 19msかかる 32
33. Fixstars Group www.fixstars.com Copyright © Fixstars Group Copyright © Fixstars Group データ転送
34. Fixstars Group www.fixstars.com Copyright © Fixstars Group データ転送 CPU-GPUデータ転送 • CPUとGPUはメモリが別 • 基本的に、CPUからGPUメモリを読み書きできないし、GPUからCPU メモリも読み書きできない • GPUで計算するには、CPUとGPUでデータを転送する必要がある • CPU-GPU間のデータ転送は、以下の方法がある • 通常のデータ転送(cudaMemcpy) • Mapped Memory • Unified Memory 34
35. Fixstars Group www.fixstars.com Copyright © Fixstars Group データ転送 通常のデータ転送(cudaMemcpy) • cudaMemcpy()で転送できる float* h_ptr = (float*)malloc(size); // Initialize input vectors ... // Allocate vectors in device memory float* d_ptr; cudaMalloc(&d_ptr , size); // Copy vectors from host memory to device memory cudaMemcpy(d_ptr , h_ptr , size, cudaMemcpyHostToDevice); … 35
36. Fixstars Group www.fixstars.com Copyright © Fixstars Group データ転送 通常のデータ転送(cudaMemcpy) • ホスト側のメモリは、できればPage-Lockedホストメモリの方が良い • Pinnedメモリとも呼ばれる • cudaHostAllocで確保する、または、mallocしたメモリをcudaHostRegisterする • 普通のmallocしたメモリと比べて、転送速度が倍くらいになる • 非同期転送(cudaMemcpyAsyncなど)する場合はこのメモリでないとダメ • 物理メモリに確保されるので、あまり多くは確保できない float* h_ptr; cudaMallocHost(&h_ptr, size); // Initialize input vectors ... // Allocate vectors in device memory float* d_ptr; cudaMalloc(&d_ptr , size); // Copy vectors from host memory to device memory cudaMemcpy(d_ptr , h_ptr , size, cudaMemcpyHostToDevice); … 36
37. Fixstars Group www.fixstars.com Copyright © Fixstars Group データ転送 Mapped Memory • ホストメモリにGPUからアクセスできるようにする機能 • cudaHostAllocまたはcudaHostRegisterで、cudaHostAllocMappedを指定すると、 GPUからもアクセスできるようになる • １度しか読み書きしないデータなら、cudaMemcpyによるデータ転送と遜色ない速 度でアクセスできるので、使っても良い • ２回以上読むようなデータは、読む度にPCIe転送が発生するので、cudaMemcpyで GPUメモリにコピーしてから使うべき 37
38. Fixstars Group www.fixstars.com Copyright © Fixstars Group データ転送 Unified Memory • Unified Memoryは、同じアドレスで、CPUからでもGPUからでも、データにア クセス可能にする機能 • cudaMallocManagedでメモリを確保する • 基本的には、アクセスしたときに、CUDAランタイムがデータを転送する • 高速化という観点からは、プログラマが明示的にデータ転送を書いた方が速い • Pascal以降でLinuxの場合は、ページ単位で転送する機能により、GPUメモリよ り多くのメモリをGPUから扱えるようになる • Pascal以前またはWindowsの場合は、cudaMallocManagedでもGPUメモリを超 える量のメモリは確保できない • Unified Virtual Address Spaceとは別の機能なので注意 • Unified Virtual Address Spaceは、GPUメモリとCPUメモリが同じ仮想アドレス スペース上に配置される機能 • アドレスから、CPUメモリかGPUメモリかを判定可能になる • cudaMemcpy*の引数cudaMemcpyKindはcudaMemcpyDefaultと書けば良い • 64bitプロセスでは常に有効 38
39. Fixstars Group www.fixstars.com Copyright © Fixstars Group データ転送 ガウシアンフィルタの実行時間で比較 • 以下の5バージョンを比較 • 通常 • 単純移植バージョン • Pinnedメモリ • cudaMemcpyをPinnedメモリで行ったバージョン • Mappedメモリ • cudaMemcpyを行わず、入出力データをMappedメモリでカーネルから直接読 み書きした場合 • 出力だけMappedメモリ • ガウシアンフィルタは入力データに複数回アクセスするので、出力データだけ、 Mappedメモリに書き込んだ場合 • Unified Memory • 入出力データのやり取りにUnified Memoryを使った場合 39
40. Fixstars Group www.fixstars.com Copyright © Fixstars Group データ転送 データ転送比較 19 8.39 9.176 6.933 0 5 10 15 20 25 30 通常 Pinnedメモリ Mapped メモリ 出力メモリ だけMapped Unified Memory データ転送も含めたガウシアンフィルタの時間 (ms) 324 計測環境 CPU: Core i7-8700 3.2GHz (6コア 12スレッド) GPU: GeForce RTX 2060 (PCIe 3.0 x16接続) 計測条件 6720x4480の画像(グレースケール 30MB) 40
41. Fixstars Group www.fixstars.com Copyright © Fixstars Group Copyright © Fixstars Group 実装方法による 性能の比較
42. Fixstars Group www.fixstars.com Copyright © Fixstars Group 実装方法による性能の比較 ループ回数を可変にしてみる __global__ void GaussianKernelArray(const uint8_t *src, uint8_t *dst, int width, int height, int step, int ks) { const float filter[5][5] = { … }; int x = blockIdx.x * blockDim.x + threadIdx.x; int y = blockIdx.y * blockDim.y + threadIdx.y; if (x < width && y < height) { float sum = 0; for (int dy = 0; dy < ks; ++dy) { for (int dx = 0; dx < ks; ++dx) { sum += filter[dy][dx] * src[(x + dx) + (y + dy) * step]; } } dst[x + y * step] = (int)(sum + 0.5f); } } ループ回数を変数で指定 カーネル 42
43. Fixstars Group www.fixstars.com Copyright © Fixstars Group 実装方法による性能の比較 ループ回数を可変にしてみる 2.605 14.36 0 2 4 6 8 10 12 14 16 ループ回数を 定数で指定 ループ回数を 変数で指定 ガウシアンフィルタ計算時間 (ms) • 5.5倍遅くなった… • 調査してみる
44. Fixstars Group www.fixstars.com Copyright © Fixstars Group 実装方法による性能の比較 PTXを見る • NVCCコンパイル時にオプションで”--keep”を付与してコンパイル • 中間生成物が残るようになる • PTXも中間生成物の１つ Visual Studioの場合 44
45. Fixstars Group www.fixstars.com Copyright © Fixstars Group 実装方法による性能の比較 PTXを見る • GaussianKernelSimple（単純移植 カーネル） • ループが完全にアンロールされて いる • フィルタの値が命令の即値になっ ている ld.global.u8 %rs6, [%rd8]; cvt.rn.f32.u16 %f11, %rs6; fma.rn.f32 %f12, %f11, 0f3C5A024A, %f10; ld.global.u8 %rs7, [%rd8+1]; cvt.rn.f32.u16 %f13, %rs7; fma.rn.f32 %f14, %f13, 0f3D744317, %f12; ld.global.u8 %rs8, [%rd8+2]; cvt.rn.f32.u16 %f15, %rs8; fma.rn.f32 %f16, %f15, 0f3DC95C2B, %f14; ld.global.u8 %rs9, [%rd8+3]; cvt.rn.f32.u16 %f17, %rs9; fma.rn.f32 %f18, %f17, 0f3D744317, %f16; ld.global.u8 %rs10, [%rd8+4]; cvt.rn.f32.u16 %f19, %rs10; fma.rn.f32 %f20, %f19, 0f3C5A024A, %f18; add.s32 %r15, %r2, 2; mad.lo.s32 %r16, %r15, %r3, %r1; cvt.s64.s32 %rd9, %r16; add.s64 %rd10, %rd3, %rd9; GaussianKernelSimple（単純移植） PTXの一部 45
46. Fixstars Group www.fixstars.com Copyright © Fixstars Group 実装方法による性能の比較 PTXを見る • GaussianKernelArray（ループ回数 可変） • 最初の方で、ローカルメモリに何 かを大量にストアしている st.local.u32 [%rd1+4], %rd12; mov.u64 %rd13, 994218967; st.local.u32 [%rd1], %rd13; st.local.u32 [%rd1+12], %rd12; mov.u64 %rd14, 1018410958; st.local.u32 [%rd1+8], %rd14; st.local.u32 [%rd1+20], %rd12; st.local.u32 [%rd1+16], %rd13; mov.u64 %rd15, 1036606507; st.local.u32 [%rd1+28], %rd15; mov.u64 %rd16, 1031029527; st.local.u32 [%rd1+24], %rd16; st.local.u32 [%rd1+36], %rd12; st.local.u32 [%rd1+32], %rd16; st.local.u32 [%rd1+44], %rd15; st.local.u32 [%rd1+40], %rd14; st.local.u32 [%rd1+52], %rd15; mov.u64 %rd17, 1042677320; st.local.u32 [%rd1+48], %rd17; GaussianKernelArray（ループ回数可変） PTXの最初の方の一部 46
47. Fixstars Group www.fixstars.com Copyright © Fixstars Group 実装方法による性能の比較 PTXを見る • GaussianKernelArray（ループ回数 可変） • アドレス計算やループカウントの 命令が多い • グローバルメモリからのロードに 加えて、ローカルメモリからもロ ードしている BB4_11: mul.lo.s64 %rd30, %rd2, 20; add.s64 %rd31, %rd1, %rd30; mul.wide.s32 %rd32, %r39, 4; add.s64 %rd33, %rd31, %rd32; add.s32 %r32, %r4, %r39; cvta.to.global.u64 %rd34, %rd9; cvt.s64.s32 %rd35, %r32; add.s64 %rd36, %rd34, %rd35; ld.global.u8 %rs3, [%rd36]; cvt.rn.f32.u16 %f20, %rs3; ld.local.f32 %f21, [%rd33]; fma.rn.f32 %f37, %f21, %f20, %f40; add.s32 %r41, %r39, 1; mov.f32 %f40, %f37; BB4_12: setp.lt.u32 %p8, %r14, 4; @%p8 bra BB4_15; GaussianKernelArray（ループ回数可変） PTXの中間あたりの一部 47
48. Fixstars Group www.fixstars.com Copyright © Fixstars Group 実装方法による性能の比較 遅くなった原因① • ローカルメモリの使用 • カーネル関数内で定義してる配列が原因 • カーネル関数内で配列を使うと、コンパイル時に参照インデックスが定数にならない 場合は、ローカルメモリに展開して解決しようとする __global__ void GaussianKernelArray(…, int ks) { const float filter[5][5] = { … }; int x = blockIdx.x * blockDim.x + threadIdx.x; int y = blockIdx.y * blockDim.y + threadIdx.y; if (x < width && y < height) { … } } これが悪い 48
49. Fixstars Group www.fixstars.com Copyright © Fixstars Group 実装方法による性能の比較 カーネルプロファイリング GaussianKernelSimple（単純移植カーネル） GaussianKernelArray（ループ回数可変） フィルタをローカルメモリに展開したせいで、メモリ使用が増大 49
50. Fixstars Group www.fixstars.com Copyright © Fixstars Group 実装方法による性能の比較 コンスタントメモリの使用 __constant__ float filter[5][5] = { … }; __global__ void GaussianKernelConstant(const uint8_t *src, uint8_t *dst, int width, int height, int step, int ks) { int x = blockIdx.x * blockDim.x + threadIdx.x; int y = blockIdx.y * blockDim.y + threadIdx.y; if (x < width && y < height) { float sum = 0; for (int dy = 0; dy < ks; ++dy) { for (int dx = 0; dx < ks; ++dx) { sum += filter[dy][dx] * src[(x + dx) + (y + dy) * step]; } } dst[x + y * step] = (int)(sum + 0.5f); } } コンスタントメモリに定義 50
51. Fixstars Group www.fixstars.com Copyright © Fixstars Group 実装方法による性能の比較 コンスタントメモリの使用 2.605 14.36 3.481 0 2 4 6 8 10 12 14 16 ループ回数を 定数で指定 ループ回数を 変数で指定 ループ回数を 変数で指定 コンスタントメモリ使用 ガウシアンフィルタ計算時間 (ms) • 速くなった • ただし、ループ回数が定 数でないので、アンロー ルできない分遅い 51
52. Fixstars Group www.fixstars.com Copyright © Fixstars Group 実装方法による性能の比較 コンスタントメモリの使用 • 速くなった GaussianKernelSimple（単純移植カーネル） GaussianKernelConstant（ループ回数可変、コンスタントメモリ使用） ループ回数定数版とほぼ同じ傾向となった が、ループをアンロールできないせいで、命令数が増え効率は落ちている 52
53. Fixstars Group www.fixstars.com Copyright © Fixstars Group 実装方法による性能の比較 Shared Memory • 入力画像に何度もアクセスするの で、Shared Memoryを使ってみる • 入力画像に5x5=25回アクセス している • 結果、速くならなかった • 実装によっては速くなるかもし れないが、L1キャッシュが効い ているので、Shared Memory を使っても効果がない場合もあ る 1.63 1.96 0 0.5 1 1.5 2 2.5 コンスタントメモリ使用 Shared Memoryを 使った ガウシアンフィルタ計算時間 (ms) 53
54. Fixstars Group www.fixstars.com Copyright © Fixstars Group 実装方法による性能の比較 L1キャッシュの使用？ • 実はすでにL1キャッシュが使われている • Volta以降はデフォルトでL1キャッシュが使われるので、特殊なことはせずと も使われる 54
55. Fixstars Group www.fixstars.com Copyright © Fixstars Group Copyright © Fixstars Group RGB画像の処理
56. Fixstars Group www.fixstars.com Copyright © Fixstars Group RGB画像の処理 • RGB画像をRGB24bitで扱うかRGBA32bitで扱うか RGB画像 RGB 24bit RGBA 32bit 性能は？ データサイズは RGB24bit < RGBA32bit だが・・・ RGB24bit RGBA32bit 56
57. Fixstars Group www.fixstars.com Copyright © Fixstars Group RGB画像の処理 ガウシアンフィルタで比較 • RGBA32bitの方が速い • 理由 • RGBA 32bitだと、4バイト を読む命令で処理できるが、 • RGB 24bitだと、1バイトず つ読んで処理するので 4.83 3.69 0 1 2 3 4 5 6 RGB24bit RGBA32bit ガウシアンフィルタ計算時間 (ms) 計測環境 CPU: Core i7-8700 3.2GHz (6コア 12スレッド) GPU: GeForce RTX 2060 OS: Windows 10 計測条件 6720x4480の画像(グレースケール)を処理 計算時間のみ、データ転送やメモリ確保などの時間を含めず 57
58. Fixstars Group www.fixstars.com Copyright © Fixstars Group RGB画像の処理 RGB 24bit のガウシアンフィルタ __global__ void GaussianKernelColor3(const uchar3 *src, uchar3 *dst, int width, int height, int step){ int x = blockIdx.x * blockDim.x + threadIdx.x; int y = blockIdx.y * blockDim.y + threadIdx.y; if (x < width && y < height) { float3 sum = { 0, 0, 0 }; for (int dy = 0; dy < 5; ++dy) { for (int dx = 0; dx < 5; ++dx) { auto s = src[(x + dx) + (y + dy) * step]; sum.x += filter[dy][dx] * s.x; sum.y += filter[dy][dx] * s.y; sum.z += filter[dy][dx] * s.z; }} uchar3 t = { (int)(sum.x + 0.5),(int)(sum.y + 0.5),(int)(sum.z + 0.5) }; dst[x + y * step] = t; }} 58
59. Fixstars Group www.fixstars.com Copyright © Fixstars Group RGB画像の処理 RGBA 32bit のガウシアンフィルタ __global__ void GaussianKernelColor4(const uchar4 *src, uchar4 *dst, int width, int height, int step){ int x = blockIdx.x * blockDim.x + threadIdx.x; int y = blockIdx.y * blockDim.y + threadIdx.y; if (x < width && y < height) { float3 sum = { 0, 0, 0 }; for (int dy = 0; dy < 5; ++dy) { for (int dx = 0; dx < 5; ++dx) { auto s = src[(x + dx) + (y + dy) * step]; sum.x += filter[dy][dx] * s.x; sum.y += filter[dy][dx] * s.y; sum.z += filter[dy][dx] * s.z; }} uchar4 t = { (int)(sum.x + 0.5),(int)(sum.y + 0.5),(int)(sum.z + 0.5),0 }; dst[x + y * step] = t; }} 59
60. Fixstars Group www.fixstars.com Copyright © Fixstars Group Thank You お問い合わせ窓口 : contact@fixstars.com