あかすくぱるふぇ

同人サークル「あかすくぱるふぇ」のブログです。

CUDAは世代間でアーキテクチャが大きく異なり、勉強していると混乱することがよくある。
そこで、基本概念に関する世代間の共通点と相違点をまとめてみた。

〇共通点
・グリッド、ブロック、スレッド
ブログラム実行に関わる論理的概念(物理的なものではない)。
スレッドは、プログラムの最小単位。
ブロックは、複数のスレッドをまとめたもの。グリッドは、複数のブロックをまとめたもの。

・WARP
32スレッドを束ねたもの(論理的概念)であり、プログラムはWARP単位で実行される。

・CUDAコア
1スレッドの処理を担当する。
ストリーミングプロセッサ(SP)とも呼ばれる。

・ストリーミングマルチプロセッサ(SM)
複数のCUDAコアをまとめたもの。1ブロックの処理を担当する。
世代によって呼び方が異なる(後述)。

〇相違点
・Fermi世代
SMを32個のCUDAコアで構成する。
GTX480は15個のSMで構成されるので、CUDAコア数は32×15=480個。
Quadro6000は14個のSMで構成されるので、CUDAコア数は32×14=448個。

・Kepler世代
SMをSMXと呼ぶ。SMXを192個のCUDAコアで構成する。
GTX780は12個のSMXで構成されるので、CUDAコア数は192×12=2304個。
QuadroK6000は15個のSMXで構成されるので、CUDAコア数は192×15=2880個。

・Maxwell世代
SMをSMMと呼ぶ。SMMを128個のCUDAコアで構成する。
GTX980は16個のSMMで構成されるので、CUDAコア数は128×16=2048個。
QuadroM6000は24個のSMMで構成されるので、CUDAコア数は128×24=3072個。

以上です。

CUDAのグローバルメモリとテクスチャメモリの速度を比較してみました。
対象とした処理は「平均値フィルタ」です。

コードは以下の通り。
// 平均値フィルタ(グローバルメモリ)
__global__ void meanFilterKernel(unsigned char *out, const unsigned char *in,
const int width, const int filterRadius, const float fscale)
{
for (int i = threadIdx.x; i < width; i += blockDim.x) {
int sum = 0;
for (int fy = -filterRadius; fy <= filterRadius; ++fy) {
int y = blockIdx.x + fy;
for (int fx = -filterRadius; fx <= filterRadius; ++fx) {
sum += in[y * width + i + fx];
}
}
out[blockIdx.x * width + i] = sum * fscale;
}
}

// 平均値フィルタ(テクスチャメモリ)
texture<unsigned char, 2> tex;
__global__ void meanFilterKernelTexture(unsigned char *out, const unsigned char *in,
const int width, const int filterRadius, const float fscale)
{
for (int i = threadIdx.x; i < width; i += blockDim.x) {
int sum = 0;
for (int y = -filterRadius; y <= filterRadius; ++y) {
for (int x = -filterRadius; x <= filterRadius; ++x) {
sum += tex2D(tex, i + x, blockIdx.x + y);
}
}
out[blockIdx.x * width + i] = sum * fscale;
}
}

void main()
{
cv::Mat img = cv::imread("Room_16000.bmp", 0);
int pixelNum = img.cols * img.rows;
int dataSize = pixelNum * sizeof(unsigned char);

unsigned char *hOut, *dIn, *dOut;
hOut = new unsigned char[pixelNum];
checkCudaErrors(cudaMalloc((unsigned char**)&dIn, dataSize));
checkCudaErrors(cudaMalloc((unsigned char**)&dOut, dataSize));
checkCudaErrors(cudaMemcpy(dIn, img.data, dataSize, cudaMemcpyHostToDevice));

int filterRadius = 5;
int filterDiameter = filterRadius * 2 + 1;
float fscale = 1.0f / filterDiameter / filterDiameter;
dim3 grid(img.rows);
dim3 block(1024);

LARGE_INTEGER freq;
QueryPerformanceFrequency(&freq);

// CPU
LARGE_INTEGER start, end;
QueryPerformanceCounter(&start);
for (int h = filterRadius; h < img.rows - filterRadius; ++h) {
for (int w = filterRadius; w < img.cols - filterRadius; ++w) {
int sum = 0;
for (int fy = -filterRadius; fy <= filterRadius; ++fy) {
int y = h + fy;
for (int fx = -filterRadius; fx <= filterRadius; ++fx) {
sum += img.data[y * img.cols + w + x];
}
}
hOut[h * img.cols + w] = sum * fscale;
}
}
QueryPerformanceCounter(&end);
 printf("CPU:%lf[ms]\n", (double)(end.QuadPart - start.QuadPart)
/ freq.QuadPart * 1000.0);

// GPU(グローバルメモリ)
QueryPerformanceCounter(&start);
meanFilterKernel <<<grid, block>>>(dOut, dIn, img.cols,
filterRadius, fscale);
checkCudaErrors(cudaDeviceSynchronize());
QueryPerformanceCounter(&end);
printf("GPU(global memory):%lf[ms]\n", (double)(end.QuadPart - start.QuadPart)
/ freq.QuadPart * 1000.0);

// GPU(テクスチャメモリ)
cudaArray *cuArray;
cudaChannelFormatDesc desc = cudaCreateChannelDesc<unsigned char>();
checkCudaErrors(cudaMallocArray(&cuArray, &desc, img.cols, img.rows));
checkCudaErrors(cudaMemcpyToArray(cuArray, 0, 0, img.data,
dataSize, cudaMemcpyHostToDevice));
checkCudaErrors(cudaBindTextureToArray(tex, cuArray));

QueryPerformanceCounter(&start);
meanFilterKernelTexture <<<grid, block>>>(dOut, dIn, img.cols,
filterRadius, fscale);
checkCudaErrors(cudaDeviceSynchronize());
QueryPerformanceCounter(&end);
printf("GPU(texture memory):%lf[ms]\n", (double)(end.QuadPart - start.QuadPart
/ freq.QuadPart * 1000.0);
}

実験は以下の環境で行いました。
CPU:i7-2600 3.4GHz
メモリ:12GB
GPU:GTX970

結果は以下のようになりました(単位は[ms]、表の左端はフィルタサイズ)。

・画像解像度2000x1000
02000

・画像解像度16000x8000
16000

想定に反してテクスチャメモリの方が遅くなってしまいました。
使い方が間違っているのかなんなのか。

CPUGPUの処理時間比は50倍程度でしょうか。
シェアードメモリを使うような処理はもっと速くなるのかもしれないですが。

☆追記
試行回数を増やして、その平均処理時間をとっても、同様の傾向でした。

Deep Neural Networkによる画風転写で沙耶の唄の世界を体験できるようにしました。
 

本記事では、体験までの道のりを紹介します。

1. 全天球画像の撮影
以下のような、(部屋の)全天球画像を撮影します。
Room_02000
Ricoh Thetaなどの360度カメラで撮影しましょう。
私はパノラマ合成で作成しましたが、撮影と合成がとても大変でした。
(壁に不自然に貼られた画像が苦労を物語っている。。。)
ただ、パノラマ合成だと、真下方向も取得できるので、その点は良いですね。

2. 画風転写
全天球画像に肉塊の画風を転写して、肉塊全天球画像(しゅごい日本語)を生成します。
style_nikuim_00
画風転写は以下の記事に載っているChainerコードを用いて行いました。
画風を変換するアルゴリズム

画像サイズが大きいとメモリ不足で実行できないようなので、全天球画像を複数領域に分割してそれぞれ画風転写し、フォトショでつなぎ合わせました。


3. UnityによるHMD体験
以下の記事の通りに、UnityでSphereに肉塊全天球画像を貼り付けます。
Unity覚え書き(全天球画像を球体に貼り付ける)

あとは、UnityのHMD用プラグインで体験できるようにするだけです。
Unity+Viveで開発する


☆感想
動画だとそれっぽく見えるんですが、HMDで体験すると今一つだったりします。
そもそも、部屋の全天球画像をHMDで見ても、スケール感が合わなくて、現実感が無いのですよね。
全天球画像ではなく、部屋のメッシュデータを取り込んでそれを使ったりすれば、現実感がグンと増すかもしれません。

↑このページのトップヘ