CUDA高速化セミナーvol.5 ～画像処理アルゴリズムの高速化２～（2022/09/29）

8.1K Views

September 29, 22

cuda 画像処理 cuda高速化高速化シリーズ cuda高速化シリーズ

スライド概要

CUDA高速化セミナーシリーズの第5回として、過去のCUDA高速化セミナーvol.1で解説したデータ転送・カーネルを書く上での基本的な注意事項を踏まえ、今回はCUDA特有の計算方法使った、バイラテラルフィルタ、転置、リダクションの実装例を紹介します。

画像処理関連分野の研究室に所属する学生や、企業のGPU搭載製品の開発部門に所属しているエンジニアにオススメの内容となっております。

＜講演内容＞
・画像処理アルゴリズムの高速化1 のおさらい
・バイラテラルフィルタ高速化
・転置
・リダクション

＜過去資料＞
・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/

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

ダウンロード

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

関連スライド

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

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

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

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

arm neon

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

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

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

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

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

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

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

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

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

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

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

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

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

各ページのテキスト

本講演の位置づけ ● CUDAに関連する様々な技術情報を、CUDA高速化セミナーとして発信しています ● vol.1 画像処理アルゴリズムの高速化 ● vol.3 ソフトウェア高速化と深層学習 ● vol.2 CUDAアーキテクチャの進化 ● vol.4 TensorRT化のワークフロー事例紹介 ● 今回は、vol.1で解説したデータ転送・カーネルを書く上での基本を踏まえ、CUDA 特有の計算方法を使った、バイラテラルフィルタ、転置、リダクションの実装例を紹介します ● こんな方に向いています ○ これから CUDA を使った画像処理をしてみたい ○ CUDA カーネルを高速化したい Copyright© Fixstars Group 4

発表者紹介冨田明彦上野晃司ソリューションカンパニーソリューション第一事業部執行役員エグゼクティブエンジニア 2008年に入社。金融、医療業界において、ソフトウェア高速化業務に携わる。その後、新規事業企画、半導体業界の事業を担当し、現職。 2016年に入社。スパコンのベンチマーク Graph500を「京」「富岳」向けに最適化し世界１位を達成。CUDAやOpenCLを使った画像処理高速化を担当。 Copyright© Fixstars Group 5

フィックスターズの強みコンピュータの性能を最大限に引き出す、ソフトウェア高速化のエキスパート集団ハードウェアの知見アルゴリズム実装力各産業・研究分野の知見目的の製品に最適なハードウェアを見抜き、その性能をフル活用するソフトウェアを開発します。ハードウェアの特徴と製品要求仕様に合わせて、アルゴリズムを改良して高速化を実現します。開発したい製品に使える技術を見抜き、実際に動作する実装までトータルにサポートします。 Copyright© Fixstars Group 7

サービス提供分野半導体産業機器金融自動車 ● NAND型フラッシュメモリ向けファームウェア開発 ● 次世代AIチップの開発環境基盤生命科学 ● Smart Factory実現への支援 ● マシンビジョンシステムの高速化 ● 自動運転の高性能化、実用化 ● ゲノム解析の高速化 ● 次世代パーソナルモビリティの研究開発 ● 医用画像処理の高速化 Copyright© Fixstars Group ● デリバティブシステムの高速化 ● HFT(アルゴリズムトレード)の高速化 ● AI画像診断システムの研究開発 9

10.

11.

画像処理アルゴリズム開発高速な画像処理需要に対して、経験豊富なエンジニアが責任を持って製品開発をご支援します。お客様の課題ご支援内容高度な画像処理や深層学習等のアルゴリズムを開発できる人材が社内に限られているアルゴリズム調査・改変課題に合ったアルゴリズム・実装手法を調査製品実装に向けて適切な改変を実施機能要件は満たせそうだが、ターゲット機器上で性能要件までクリアできるか不安深層学習ネットワーク精度の改善様々な手法を駆使して深層学習ネットワークの精度を改善製品化に結びつくような研究ができていない論文調査・改善活動論文調査から最先端の手法の探索性能向上に向けた改善活動を継続 Copyright© Fixstars Group 11

12.

GPU向け高速化高性能なGPUの本来の性能を十分に引き出し、ソフトウェアの高速化を実現します。お客様の課題ご支援内容 GPUで計算してみたが期待した性能が出ない GPU高速化に関するコンサルティング GPU/CPUを組み合わせた全体として最適な設 CPU・GPU混在環境でのシステム設計計がしたいアルゴリズムのGPU向け移植原価を維持したまま機能を追加するため、もう少し処理を速くしたい GPUプログラム高速化品質確保のため、精度を上げたく演算量は増継続的な精度向上えるが性能は維持したい Copyright© Fixstars Group 12

13.

AI・深層学習向け技術支援 AIを使うためのハードウェア選定や、高速な計算を実現するソフトウェア開発技術で、お客様の製品開発を支援します。お客様の課題ご支援内容推論精度を維持したまま計算時間を短縮したい組込みデバイス向けにAIモデルを軽量化 AIモデル設計データの前処理・後処理したい推論精度の改善学習計算を高速化して研究開発を効率化したい分散処理による学習高速化精度と計算時間を両立するAIモデルをモデル圧縮・推論の高速化開発したい Copyright© Fixstars Group 13

14.

15.

16.

17.

ガウシアンフィルタ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 17

18.

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

19.

20.

21.

[beta]

__global__ void BilateralKernelSimple( const uint8_t *src, uint8_t *dst, int width, int
height, int step, float sigma)
{
必ず”f”を付ける。付けないと doubleの演
int x = blockIdx.x * blockDim.x + threadIdx.x;
算になって、Tesla以外では、
int y = blockIdx.y * blockDim.y + threadIdx.y;
かなり遅くなるので注意
if (x < width && y < height) {
float coef = 1.0 / sqrtf(2 * 3.1415926f * sigma * sigma);
float coef2 = -1.0 / (2 * sigma * sigma);
float c_sum = 0;
float f_sum = 0;
int val0 = src[x + y * step];
for (int dy = 0; dy < 3; ++dy) {
for (int dx = 0; dx < 3; ++dx) {
int val = src[(x + dx) + (y + dy) * step];
int diff = val - val0;
float w = filter3[dy][dx] * coef * expf(diff * diff * coef2);
f_sum += w;
重みの計算
c_sum += w * val;
}
}
dst[x + y * step] = (int)(c_sum / f_sum + 0.5f);
Copyright© Fixstars Group
}}

21

22.

23.

[beta]

__global__ void BilateralKernelSimple( const uint8_t *src, uint8_t *dst, int width, int
height, int step, float sigma)
{
int x = blockIdx.x * blockDim.x + threadIdx.x;
int y = blockIdx.y * blockDim.y + threadIdx.y;
割り算、sqrt
if (x < width && y < height) {
float coef = 1.0 / sqrtf(2 * 3.1415926f * sigma * sigma);
float coef2 = -1.0 / (2 * sigma * sigma);
float c_sum = 0;
float f_sum = 0;
割り算
int val0 = src[x + y * step];
for (int dy = 0; dy < 3; ++dy) {
for (int dx = 0; dx < 3; ++dx) {
int val = src[(x + dx) + (y + dy) * step];
int diff = val - val0;
float w = filter3[dy][dx] * coef * expf(diff * diff * coef2);
f_sum += w;
c_sum += w * val;
exp
割り算
}
}
dst[x + y * step] = (int)(c_sum / f_sum + 0.5f);
Copyright© Fixstars Group
}}

重い演算が多い

23

24.

[beta]

__global__ void BilateralKernelFast( const uint8_t *src, uint8_t *dst, int width, int height,
int step, float sigma)
{
int x = blockIdx.x * blockDim.x + threadIdx.x;
int y = blockIdx.y * blockDim.y + threadIdx.y;
1/sqrtf(x) のintrinsic
if (x < width && y < height) {
float coef = __frsqrt_rn(2 * 3.1415926f * sigma * sigma);
float coef2 = __frcp_rn(-2 * sigma * sigma);
float c_sum = 0;
1/x のintrinsic
float f_sum = 0;
int val0 = src[x + y * step];
for (int dy = 0; dy < 3; ++dy) {
for (int dx = 0; dx < 3; ++dx) {
int val = src[(x + dx) + (y + dy) * step];
int diff = val - val0;
float w = filter3[dy][dx] * coef * __expf(diff * diff * coef2);
f_sum += w;
c_sum += w * val;
expf(x) の高速版
}
x/y の高速版
}
dst[x + y * step] = (int)(__fdividef(c_sum, f_sum) + 0.5f);
Copyright© Fixstars Group
}}

24

25.

https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#intrinsic-functions

26.

27.

28.

29.

転置単純に書いてみる __global__ width, int { int x int y if (x void TransposeKernelSimple( const uint8_t *src, uint8_t *dst, int height) = blockIdx.x * blockDim.x + threadIdx.x; = blockIdx.y * blockDim.y + threadIdx.y; < width && y < height) dst[y + x * height] = src[x + y * width]; } xとyを逆にするだけ Copyright© Fixstars Group 29

30.

31.

転置単純に書いてみる __global__ width, int { int x int y if (x void TransposeKernelSimple( const uint8_t *src, uint8_t *dst, int height) = blockIdx.x * blockDim.x + threadIdx.x; = blockIdx.y * blockDim.y + threadIdx.y; < width && y < height) dst[y + x * height] = src[x + y * width]; } 書き込みが全くコアレスアクセスになっていない Copyright© Fixstars Group 31

32.

33.

34.

転置 Shared Memoryを使う __global__ void TransposeKernelShared( const uint8_t *src, uint8_t *dst, int width, int height) { int tx = threadIdx.x; int ty = threadIdx.y; int xbase = blockIdx.x * blockDim.x; int ybase = blockIdx.y * blockDim.y; __shared__ uint8_t sbuf[16][16]; { int x = xbase + tx; 一旦Shared Memoryに格納 int y = ybase + ty; if (x < width && y < height) sbuf[ty][tx] = src[x + y * width]; } __syncthreads(); { int x = xbase + ty; 書き込みもコアレスアクセスで int y = ybase + tx; できるようにする if (x < width && y < height) dst[y + x * height] = sbuf[tx][ty]; }} Copyright© Fixstars Group 34

35.

36.

37.

38.

転置バンクコンフリクト回避 __global__ void TransposeKernelFast( const uint8_t *src, uint8_t *dst, int width, int height){ int tx = threadIdx.x; int ty = threadIdx.y; int xbase = blockIdx.x * blockDim.x; パディングを追加 int ybase = blockIdx.y * blockDim.y; Shared Memoryのバンクは __shared__ uint8_t sbuf[16][16+4]; 4バイトインターリーブされているので、 { int x = xbase + tx; 4バイトパディングを追加する int y = ybase + ty; if (x < width && y < height) sbuf[ty][tx] = src[x + y * width]; } __syncthreads(); { int x = xbase + ty; int y = ybase + tx; if (x < width && y < height) dst[y + x * height] = sbuf[tx][ty]; }} Copyright© Fixstars Group 38

39.

40.

41.

42.

転置 1スレッドあたり処理量を増やす __global__ void TransposeKernelFast2( const uint8_t *src, uint8_t *dst, int width, int height){ int tx = threadIdx.x; int ty = threadIdx.y; int xbase = blockIdx.x * 32; １スレッドが4要素処理するように修正 int ybase = blockIdx.y * 32; __shared__ uint8_t sbuf[32][32+4]; { int x = xbase + tx; if (x < width) { int yend = min(ybase + 32, height); for (int tyy = ty, y = ybase + ty; y < yend; tyy += 8, y += 8) { sbuf[tyy][tx] = src[x + y * width]; }}} __syncthreads(); { int y = ybase + tx; if (y < height) { int xend = min(xbase + 32, width); for (int tyy = ty, x = xbase + ty; x < xend; tyy += 8, x += 8) { dst[y + x * height] = sbuf[tx][tyy]; Copyright© Fixstars Group }}}} 42

43.

44.

45.

46.

リダクション Y軸リダクション • 1スレッド１列担当 • コアレスアクセスになっていることに注意 __global__ void ReduceHKernelSimple( const uint8_t *src, float *dst, int width, int height) { int x = blockIdx.x * blockDim.x + threadIdx.x; if (x < width) { float sum = 0; for (int y = 0; y < height; ++y) { sum += src[x + y * width]; } dst[x] = sum; } } Copyright© Fixstars Group Y軸リダクション 46

47.

48.

リダクション Y軸リダクション • 列を分割して並列数を増やす ○ 1列1スレッド→ceil(行数/128)スレッド __global__ void ReduceHKernelFast( const uint8_t *src, float *dst, int width, int height) { int x = blockIdx.x * blockDim.x + threadIdx.x; 128行ごとに分割して処理する int y = blockIdx.y * 128; if (x < width) { float sum = 0; for (int yend = min(y + 128, height); y < yend; ++y) { sum += src[x + y * width]; } atomicAdd(&dst[x], sum); dstはこのカーネルを呼び出す前に } ゼロ初期化しておく } Copyright© Fixstars Group 48

49.

50.

リダクション X軸リダクション • Y軸リダクションと同じように実装 __global__ void ReduceWKernelSimple( const uint8_t *src, float *dst, int width, int height) { int y = blockIdx.x * blockDim.x + threadIdx.x; int x = blockIdx.y * 128; X軸リダクション if (y < height) { float sum = 0; for (int xend = min(x + 128, width); x < xend; ++x) { sum += src[x + y * width]; } atomicAdd(&dst[y], sum); } } Copyright© Fixstars Group 50

51.

52.

リダクション X軸リダクション __global__ width, int { int y int x void ReduceWKernelSimple( const uint8_t *src, float *dst, int height) = blockIdx.x * blockDim.x + threadIdx.x; = blockIdx.y * 128; if (y < height) { float sum = 0; for (int xend = min(x + 128, width); x < xend; ++x) { sum += src[x + y * width]; } このアクセスが全く atomicAdd(&dst[y], sum); コアレスアクセスでない } } Copyright© Fixstars Group 52

53.

54.

55.

リダクションパラレルリダクション • 1行を1ブロックが担当 __global__ void ReduceWKernelFast( const uint8_t *src, float *dst, int width, int height) { 1ブロック512スレッドで int tid = threadIdx.x; コードを書いた場合 int y = blockIdx.y; __shared__ float sbuf[512]; float sum = 0; for (int x = tid; x < width; x += 512) { sum += src[x + y * width]; 512要素までのリダクションは普通にス } レッドごとに計算 sbuf[tid] = sum; __syncthreads(); Shared Memoryに書いて sum = ReduceFunc(tid, sbuf); パラレルリダクションを呼び出す if (tid == 0) dst[y] = sum; } Copyright© Fixstars Group 55

56.

リダクションパラレルリダクション __device__ float ReduceFunc( int tid, float* buf) { if (tid < 256) { buf[tid] += buf[tid + 256]; } __syncthreads(); if (tid < 128) { buf[tid] += buf[tid + 128]; } __syncthreads(); if (tid < 64) { buf[tid] += buf[tid + 64]; } __syncthreads(); float sum; if (tid < 32) { 32スレッドまでは sum = buf[tid] + buf[tid + 32]; __syncthreads()を使って計算 sum += __shfl_down_sync(0xffffffff, sum, 16); sum += __shfl_down_sync(0xffffffff, sum, 8); sum += __shfl_down_sync(0xffffffff, sum, 4); sum += __shfl_down_sync(0xffffffff, sum, 2); 32スレッドになったら、Warp sum += __shfl_down_sync(0xffffffff, sum, 1); Shuffleで計算 } return sum; } Copyright© Fixstars Group 56

57.

58.

本セミナーのまとめ ● バイラテラルフィルタ高速化 CUDA組み込み関数を使って演算を軽量化 ○ ● ● 転置 ○ Shared Memoryを使ったメモリアクセス最適化 ○ バンクコンフリクト回避リダクション ○ X軸、Y軸方向のリダクション ○ メモリのアクセス方向を意識した計算 ○ 水平方向のリダクションを高速に行うパラレルリダクション Copyright© Fixstars Group 58

59.

CUDA高速化セミナーvol.5 ～画像処理アルゴリズムの高速化２～（2022/09/29）

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

関連スライド

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

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

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

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

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

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

各ページのテキスト

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