17.9K Views
September 29, 22
スライド概要
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 画像処理アルゴリズムの高速化2:https://www.docswell.com/s/fixstars/ZQ81QX-20220929
フィックスターズは、コンピュータの性能を最大限に引き出すソフトウェア開発のスペシャリストです。車載、産業機器、金融、医療など、幅広い分野での開発経験があります。また、ディープラーニングや機械学習などの最先端技術にも力を入れています。 並列化や最適化技術を駆使して、マルチコアCPU、GPU、FPGA、量子アニーリングマシンなど、さまざまなハードウェアでソフトウェアを高速化するサービスを提供しています。さらに、長年の経験から培ったハードウェアの知識と最適化ノウハウを活かし、高精度で高性能なアルゴリズムの開発も行っています。 ・開催セミナー一覧:https://www.fixstars.com/ja/seminar ・技術ブログ :https://proc-cpuinfo.fixstars.com/
CUDA高速化セミナー vol.5 画像処理アルゴリズムの高速化2 Copyright© Fixstars Group
本日のAgenda ● はじめに ● フィックスターズの紹介 ● 画像処理アルゴリズムの高速化1 のおさらい ● バイラテラルフィルタ高速化 ● 転置 ● リダクション ● まとめ Copyright© Fixstars Group 2
はじめに Copyright© Fixstars Group
本講演の位置づけ ● CUDAに関連する様々な技術情報を、CUDA高速化セミナーとして発信しています ● vol.1 画像処理アルゴリズムの高速化 ● vol.3 ソフトウェア高速化と深層学習 ● vol.2 CUDAアーキテクチャの進化 ● vol.4 TensorRT化のワークフロー事例紹介 ● 今回は、vol.1で解説したデータ転送・カーネルを書く上での基本を踏まえ、CUDA 特有の 計算方法を使った、バイラテラルフィルタ、転置、リダクションの実装例を紹介します ● こんな方に向いています ○ これから CUDA を使った画像処理をしてみたい ○ CUDA カーネルを高速化したい Copyright© Fixstars Group 4
発表者紹介 冨田 明彦 上野 晃司 ソリューションカンパニー ソリューション第一事業部 執行役員 エグゼクティブエンジニア 2008年に入社。金融、医療業界において、 ソフトウェア高速化業務に携わる。その後、 新規事業企画、半導体業界の事業を担当し、 現職。 2016年に入社。スパコンのベンチマーク Graph500を「京」「富岳」向けに最適化し 世界1位を達成。CUDAやOpenCLを使った 画像処理高速化を担当。 Copyright© Fixstars Group 5
フィックスターズの ご紹介 Copyright© Fixstars Group
フィックスターズの強み コンピュータの性能を最大限に引き出す、ソフトウェア高速化のエキスパート集団 ハードウェアの知見 アルゴリズム実装力 各産業・研究分野の知見 目的の製品に最適なハードウェアを見抜 き、その性能をフル活用するソフトウェア を開発します。 ハードウェアの特徴と製品要求仕様に合わ せて、アルゴリズムを改良して高速化を実 現します。 開発したい製品に使える技術を見抜き、実 際に動作する実装までトータルにサポート します。 Copyright© Fixstars Group 7
サービス概要 お客様専任のエンジニアが直接ヒアリングを行い、高速化を実現するために乗り越えるべき 課題や問題を明確にしていきます。 高速化のワークフロー コンサルティング 高速化 サポート 先行技術調査 アルゴリズムの改良・開発 レポートやコードへのQ&A 性能評価・ボトルネックの特定 ハードウェアへの最適化 実製品への組込み支援 レポート作成 Copyright© Fixstars Group 8
サービス提供分野 半導体 産業機器 金融 自動車 ● NAND型フラッシュメモリ向け ファームウェア開発 ● 次世代AIチップの開発環境基盤 生命科学 ● Smart Factory実現への支援 ● マシンビジョンシステムの高速化 ● 自動運転の高性能化、実用化 ● ゲノム解析の高速化 ● 次世代パーソナルモビリティの 研究開発 ● 医用画像処理の高速化 Copyright© Fixstars Group ● デリバティブシステムの高速化 ● HFT(アルゴリズムトレード)の高速化 ● AI画像診断システムの研究開発 9
サービス領域 様々な領域でソフトウェア高速化サービスを提供しています。大量データの高速処理は、 お客様の製品競争力の源泉となっています。 組込み高速化 GPU向け高速化 AI・深層学習 画像処理・アルゴリズム 開発 FPGAを活用した システム開発 分散並列システム開発 量子コンピューティング 自動車向け フラッシュメモリ向け ソフトウェア開発 ファームウェア開発 Copyright© Fixstars Group 10
画像処理アルゴリズム開発 高速な画像処理需要に対して、経験豊富なエンジニアが 責任を持って製品開発をご支援します。 お客様の課題 ご支援内容 高度な画像処理や深層学習等のアルゴリズム を開発できる人材が社内に限られている アルゴリズム調査・改変 課題に合ったアルゴリズム・実装手法を調査 製品実装に向けて適切な改変を実施 機能要件は満たせそうだが、ターゲット機器 上で性能要件までクリアできるか不安 深層学習ネットワーク精度の改善 様々な手法を駆使して深層学習ネットワークの精度を改善 製品化に結びつくような研究ができていない 論文調査・改善活動 論文調査から最先端の手法の探索 性能向上に向けた改善活動を継続 Copyright© Fixstars Group 11
GPU向け高速化 高性能なGPUの本来の性能を十分に引き出し、 ソフトウェアの高速化を実現します。 お客様の課題 ご支援内容 GPUで計算してみたが期待した性能が出ない GPU高速化に関するコンサルティング GPU/CPUを組み合わせた全体として最適な設 CPU・GPU混在環境でのシステム設計 計がしたい アルゴリズムのGPU向け移植 原価を維持したまま機能を追加するため、も う少し処理を速くしたい GPUプログラム高速化 品質確保のため、精度を上げたく演算量は増 継続的な精度向上 えるが性能は維持したい Copyright© Fixstars Group 12
AI・深層学習向け技術支援 AIを使うためのハードウェア選定や、高速な計算を実現する ソフトウェア開発技術で、お客様の製品開発を支援します。 お客様の課題 ご支援内容 推論精度を維持したまま計算時間を短縮 したい 組込みデバイス向けにAIモデルを軽量化 AIモデル設計 データの前処理・後処理 したい 推論精度の改善 学習計算を高速化して研究開発を効率化 したい 分散処理による学習高速化 精度と計算時間を両立するAIモデルを モデル圧縮・推論の高速化 開発したい Copyright© Fixstars Group 13
前回の復習 Copyright© Fixstars Group
NVIDIA Nsight Compute • カーネルプロファイラをサポート Copyright© Fixstars Group 15
ガウシアンフィルタ ⊗ カーネル Copyright© Fixstars Group 16
ガウシアンフィルタ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 17
本日説明するコード ● ↓ここにあります ● https://github.com/fixstars/CudaOptimizeSample/blob/master/CudaO ptimizeSample/kernel.cu ● (前回と同じファイルです) Copyright© Fixstars Group 18
バイラテラルフィルタ 高速化 Copyright© Fixstars Group
バイラテラルフィルタ高速化 • ガウシアンフィルタと同じように、画像と重みを掛け合わせるが、重みが画 素値の差によって変わるようにしたもの Copyright© Fixstars Group 20
__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
バイラテラルフィルタ高速化 カーネルプロファイリング • かなり演算ネック Copyright© Fixstars Group 22
__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
__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
バイラテラルフィルタ高速化 高速版の誤差 ulp = Unit in the last place https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#intrinsic-functions Copyright© Fixstars Group 25
バイラテラルフィルタ高速化 最適化前 最適化後 バイラテラルフィルタ 計算時間 (ms) 最適化前 2.48 最適化後Copyright© Fixstars Group 1.63 26
転置 Copyright© Fixstars Group
転置 Copyright© Fixstars Group 28
転置 単純に書いてみる __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
転置 カーネルプロファイル メモリがネック Loadに比べてStoreの L2⇔DRAM Transactionが かなり多い Copyright© Fixstars Group 30
転置 単純に書いてみる __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
転置 単純に書いてみる メモリ連続方向 Warpのアクセスするメモリ メモリ連続方向 Store Load NG! 完全に飛び飛びのアクセス コアレスアクセスになっている Copyright© Fixstars Group 32
転置 Shared Memoryを使う Warpのアクセスするメモリ メモリ連続方向 メモリ連続方向 Shared Memory Shared Memory 同じメモリ 書き込みもコアレスアクセス でできるようにする Copyright© Fixstars Group 33
転置 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
転置 Shared Memoryを使う メモリネックでなくなった LoadとStoreで L2⇔DRAM Transactionが 同じになった Copyright© Fixstars Group 35
転置 Shared Memoryを使う Copyright© Fixstars Group 計測環境 CPU: Core i7-8700 3.2GHz (6コア 12スレッド) GPU: GeForce RTX 2060 計測条件 6720x4480の画像(グレースケール)を処理 計算時間のみで、データ転送やメモリ確保などの 時間を含めず 36
転置 Shared Memoryを使う • ただし、Shared Memory のバンクコンフリクトが発生している Copyright© Fixstars Group 37
転置 バンクコンフリクト回避 __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
転置 バンクコンフリクト回避 • バンクコンフリクトがゼロになった Copyright© Fixstars Group 39
転置 バンクコンフリクト回避 Copyright© Fixstars Group 計測環境 CPU: Core i7-8700 3.2GHz (6コア 12スレッド) GPU: GeForce RTX 2060 計測条件 6720x4480の画像(グレースケール)を処理 計算時間のみで、データ転送やメモリ確保などの 時間を含めず 40
転置 1スレッドあたり処理量を増やす • 1スレッドあたりの処理量が少なすぎる ○ • カーネルが32命令しかない 1スレッドが処理する要素数を増やす Copyright© Fixstars Group カーネルのアセンブラ
転置 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; 1スレッドが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
転置 1スレッドあたり処理量を増やす Copyright© Fixstars Group 計測環境 CPU: Core i7-8700 3.2GHz (6コア 12スレッド) GPU: GeForce RTX 2060 計測条件 6720x4480の画像(グレースケール)を処理 計算時間のみで、データ転送やメモリ確保などの 時間を含めず 43
リダクション Copyright© Fixstars Group
リダクション • 2次元データのリダクション • 2つの軸のリダクションをそれぞれ実装してみる x Y軸リダクション X軸リダクション メモリ連続方向 メモリ連続方向 y Copyright© Fixstars Group 45
リダクション Y軸リダクション • 1スレッド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
リダクション Y軸リダクション • 6720x4480だとブロック数が7個しかない ○ 1スレッド1列担当なので ceil(6720/1024)=7 Copyright© Fixstars Group 47
リダクション 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
リダクション Y軸リダクション 計測環境 CPU: Core i7-8700 3.2GHz (6コア 12スレッド) GPU: GeForce RTX 2060 計測条件 6720x4480の画像(グレースケール)を処理 計算時間のみで、データ転送やメモリ確保などの 時間を含めず Copyright© Fixstars Group 49
リダクション 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
リダクション X軸リダクション • Y軸リダクションと同じように実装 • 2.39ms …遅い Y軸と同じ方法 X軸リダクション時間 (ms) 2.39 X軸リダクション 参考 1列1スレッド 列も分割 Y軸リダクション時間 (ms) 0.28 0.115 Copyright© Fixstars Group 51
リダクション 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
リダクション X軸リダクション Warpのアクセスするメモリ X軸リダクション このアクセスはダメ こうアクセスしたい! Copyright© Fixstars Group 水平方向の リダクション が必要 53
リダクション 水平方向のリダクションが必要 パラレルリダクション Copyright© Fixstars Group 54
リダクション パラレルリダクション • 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
リダクション パラレルリダクション __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
リダクション パラレルリダクション Copyright© Fixstars Group 57
本セミナーのまとめ ● バイラテラルフィルタ高速化 CUDA組み込み関数を使って演算を軽量化 ○ ● ● 転置 ○ Shared Memoryを使ったメモリアクセス最適化 ○ バンクコンフリクト回避 リダクション ○ X軸、Y軸方向のリダクション ○ メモリのアクセス方向を意識した計算 ○ 水平方向のリダクションを高速に行うパラレルリダクション Copyright© Fixstars Group 58
Thank you! お問い合わせ窓口 : [email protected] Copyright © Fixstars Group