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