CUDA高速化セミナー vol.1 －画像処理アルゴリズムの高速化－（2022/05/27）

19.4K Views

May 27, 22

#cuda #gpu #画像処理 #cuda高速化 #高速化シリーズ #cuda高速化シリーズ #CUDA #GPU高速化 #アルゴリズム #Fixstars

スライド概要

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 　画像処理アルゴリズムの高速化２：https://www.docswell.com/s/fixstars/ZQ81QX-20220929

株式会社フィックスターズ

@fixstars

スライド一覧

フィックスターズは、コンピュータの性能を最大限に引き出すソフトウェア開発のスペシャリストです。車載、産業機器、金融、医療など、幅広い分野での開発経験があります。また、ディープラーニングや機械学習などの最先端技術にも力を入れています。並列化や最適化技術を駆使して、マルチコアCPU、GPU、FPGA、量子アニーリングマシンなど、さまざまなハードウェアでソフトウェアを高速化するサービスを提供しています。さらに、長年の経験から培ったハードウェアの知識と最適化ノウハウを活かし、高精度で高性能なアルゴリズムの開発も行っています。　　　　　　・開催セミナー一覧：https://www.fixstars.com/ja/seminar　　・技術ブログ：https://proc-cpuinfo.fixstars.com/

またはPlayer版

埋め込む »CMSなどでJSが使えない場合

ダウンロード

ダウンロード(pdf - 2.26MB)

関連スライド

Kaggleスコアアップセミナー～画像系コンペ入門[前編]（2023/08/02）

kaggle 画像処理機械学習深層学習ディープラーニングコンペ初心者 kaggleスコアアップシリーズ

株式会社フィックスターズ 63.5K

いまさら聞けない ARM を使ったNEONの基礎と活用事例（2021/08/05）

arm neon

株式会社フィックスターズ 38.6K

ROS2自律走行実現に向けて 2 ～次世代ロボット開発フレームワークROS2のビルドシステム徹底理解～（2022/11/30）

ros2 gpu ロボット自律走行 ros2シリーズ高速化シリーズ

株式会社フィックスターズ 30.6K

Kaggleスコアアップセミナー～画像系コンペ入門[後編]（2023/09/26）

kaggle 画像処理機械学習ディープラーニングコンペ初心者 kaggleスコアアップシリーズ

株式会社フィックスターズ 29.9K

いまさら聞けない！CUDA高速化入門～プログラミングモデルとアーキテクチャの解説、高速化の実践～（2021/10/29）

cuda gpu cuda高速化高速化シリーズ

株式会社フィックスターズ 28.9K

CUDA高速化セミナー vol.2 ～CUDAアーキテクチャの進化～（2022/06/23）

gpu cuda cuda高速化高速化シリーズ cuda高速化シリーズ

株式会社フィックスターズ 27K

各ページのテキスト

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 ソフトウェア高速化サービス (概要) お客様のソースコードをご提供いただき、最適化やアルゴリズムの改良を施して高速化してお返ししますオリジナルソースコードのご提供当社コンサルティング高速化したソースコード高速化お客様サポート要件分析アルゴリズムの改良・開発実製品への組込み支援先行研究等の調査ハードウェアへの最適化レポートやコードへの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

10.

Fixstars Group www.fixstars.com 画像処理・アルゴリズム開発サービス • お客様の課題 • 高度な画像処理や深層学習等のアルゴリズム開発を行える人材が社内に限られている • 考案中のアルゴリズムで機能要件は満たせそうだが、ターゲット機器上で性能要件までクリアできるか不安 • 製品化に結びつくような研究ができていない • 弊社の支援内容 • 課題に応じたアルゴリズム調査 • 深層学習ネットワーク精度改善、推論高速化手法調査 • 論文調査、実装出展：https://www.cs.toronto.edu/~frossard/post/vgg16/ Copyright © Fixstars Group 10

https://www.cs.toronto.edu/~frossard/post/vgg16/

11.

Fixstars Group www.fixstars.com AI・深層学習関連サービス • ディープラーニングの包括的開発技術 • ネットワーク設計からターゲットデバイスでの高速化のノウハウ • 大規模システムからエッジコンピューティングまでの開発実績ネットワーク設計データの前処理、データ拡張精度改善分散処理による学習高速化各種DLフレームワークターゲットデバイスへのポーティング及び推論高速化 ■ ARM, GPU, DSP ■ SIMD,NEON,CUDA,TensorRT モデル圧縮 - 量子化 - 枝刈り - 蒸留クラウド・サーバエッジ Copyright © Fixstars Group 11

12.

Fixstars Group www.fixstars.com GPU向け高速化サービス • お客様の課題 • • GPU 高速化の知見がない自力で GPU に乗せてみたものの望む性能が出ない • 弊社の支援内容 • • • GPU 高速化に関するコンサルティングボトルネック調査、GPU プログラムの高速化 CPU/GPU が混在するヘテロジニアス環境での最適化 10～150 倍の高速化事例あり 12 Copyright © Fixstars Group

13.

14.

15.

16.

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

17.

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

18.

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

19.

20.

21.

22.

Fixstars Group www.fixstars.com プロファイラー • プロファイラーは性能を分析するツール • • VoltaまでのGPUなら • • CUDAのボトルネック解析や最適化に必須 NVIDIA Visual Profiler Turing世代以降のGPUの場合 • • NVIDIA Nsight Systems NVIDIA Nsight Compute 22 Copyright © Fixstars Group

23.

24.

25.

26.

https://github.com/fixstars/CudaOptimizeSample/blob/master/CudaOptimizeSample/kernel.cu

27.

28.

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

29.

Fixstars Group www.fixstars.com ガウシアンフィルタCUDA化スレッド割り当て • 1スレッドが出力1ピクセルを担当 • ブロックの最大スレッド数は1024なので、１ブロック 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

30.

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

31.

[beta]

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
}

32.

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

33.

34.

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

35.

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

36.

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

37.

Fixstars Group www.fixstars.com データ転送 Mapped Memory • ホストメモリにGPUからアクセスできるようにする機能 • cudaHostAllocまたはcudaHostRegisterで、cudaHostAllocMappedを指定すると、 GPUからもアクセスできるようになる • １度しか読み書きしないデータなら、cudaMemcpyによるデータ転送と遜色ない速度でアクセスできるので、使っても良い • ２回以上読むようなデータは、読む度にPCIe転送が発生するので、cudaMemcpyで GPUメモリにコピーしてから使うべき 37 Copyright © Fixstars Group

38.

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

39.

Fixstars Group www.fixstars.com データ転送ガウシアンフィルタの実行時間で比較 • 以下の5バージョンを比較 • 通常 • • Pinnedメモリ • • cudaMemcpyを行わず、入出力データをMappedメモリでカーネルから直接読み書きした場合出力だけMappedメモリ • • cudaMemcpyをPinnedメモリで行ったバージョン Mappedメモリ • • 単純移植バージョンガウシアンフィルタは入力データに複数回アクセスするので、出力データだけ、 Mappedメモリに書き込んだ場合 Unified Memory • 入出力データのやり取りにUnified Memoryを使った場合 39 Copyright © Fixstars Group

40.

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

41.

42.

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

43.

44.

45.

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

46.

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

47.

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

48.

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

49.

50.

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

51.

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

52.

Fixstars Group www.fixstars.com 実装方法による性能の比較コンスタントメモリの使用 GaussianKernelSimple（単純移植カーネル） • 速くなった GaussianKernelConstant（ループ回数可変、コンスタントメモリ使用）ループ回数定数版とほぼ同じ傾向となったが、ループをアンロールできないせいで、命令数が増え効率は落ちている Copyright © Fixstars Group 52

53.

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 使った

54.

55.

56.

57.

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

58.

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

59.

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

60.